Skip to content

Commit 868d57c

Browse files
authored
Add HIP test for ray-tracing workload (#74)
1 parent 376a3a3 commit 868d57c

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

61 files changed

+3112647
-3
lines changed

External/HIP/CMakeLists.txt

+21-3
Original file line numberDiff line numberDiff line change
@@ -12,8 +12,26 @@ macro(create_local_hip_tests VariantSuffix)
1212
list(APPEND HIP_LOCAL_TESTS empty)
1313
list(APPEND HIP_LOCAL_TESTS with-fopenmp)
1414
list(APPEND HIP_LOCAL_TESTS saxpy)
15+
list(APPEND HIP_LOCAL_TESTS InOneWeekend)
16+
list(APPEND HIP_LOCAL_TESTS TheNextWeek)
17+
18+
# Copy files needed for ray-tracing tests.
19+
file(GLOB IMAGE_FILES "workload/ray-tracing/images/*.jpg" "workload/ray-tracing/images/*.png")
20+
file(COPY ${IMAGE_FILES} DESTINATION "${CMAKE_CURRENT_BINARY_DIR}")
21+
1522
foreach(_hip_test IN LISTS HIP_LOCAL_TESTS)
16-
create_one_local_test(${_hip_test} ${_hip_test}.hip
23+
set(test_source "${_hip_test}.hip")
24+
25+
if(_hip_test STREQUAL "TheNextWeek" OR _hip_test STREQUAL "InOneWeekend")
26+
file(GLOB REF_PPM_FILES "workload/ray-tracing/${_hip_test}/*.ppm")
27+
file(COPY ${REF_PPM_FILES} DESTINATION "${CMAKE_CURRENT_BINARY_DIR}")
28+
set(test_source "workload/ray-tracing/${_hip_test}/main.cc")
29+
# need -mfma to enable FMA in host code
30+
set_source_files_properties(${test_source} PROPERTIES
31+
COMPILE_FLAGS "-xhip -mfma")
32+
endif()
33+
34+
create_one_local_test(${_hip_test} ${test_source}
1735
${VariantOffload} ${VariantSuffix}
1836
"${VariantCPPFLAGS}" "${VariantLibs}")
1937
endforeach()
@@ -64,15 +82,15 @@ macro(create_hip_tests)
6482
COMMENT "Run all simple HIP tests")
6583

6684
if(NOT AMDGPU_ARCHS)
67-
list(APPEND AMDGPU_ARCHS "gfx906;gfx90a;gfx1030;gfx1100")
85+
list(APPEND AMDGPU_ARCHS "gfx906;gfx90a;gfx1030;gfx1100;native")
6886
endif()
6987

7088
foreach(_RocmPath ${ROCM_PATHS})
7189
get_version(_RocmVersion ${_RocmPath})
7290
set(_HIP_Suffix "hip-${_RocmVersion}")
7391
# Set up HIP test flags
7492
set(_HIP_CPPFLAGS --rocm-path=${_RocmPath})
75-
set(_HIP_LDFLAGS --rocm-path=${_RocmPath} --hip-link -rtlib=compiler-rt -frtlib-add-rpath)
93+
set(_HIP_LDFLAGS --rocm-path=${_RocmPath} --hip-link -rtlib=compiler-rt -unwindlib=libgcc -frtlib-add-rpath)
7694

7795
# Unset these for each iteration of rocm path.
7896
set(_ArchFlags)
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,6 @@
1+
image width = 1200 height = 675
2+
block size = (16, 16) grid size = (75, 43)
3+
Start rendering by GPU.
4+
Done.
5+
gpu.ppm and ref.ppm are the same.
6+
exit 0
+49
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,49 @@
1+
Running quads
2+
image width = 400 height = 400
3+
block size = (16, 16) grid size = (25, 25)
4+
Start rendering by GPU.
5+
Done.
6+
quads_gpu.ppm and quads_ref.ppm are the same.
7+
Running earth
8+
image width = 400 height = 225
9+
block size = (16, 16) grid size = (25, 15)
10+
Start rendering by GPU.
11+
Done.
12+
earth_gpu.ppm and earth_ref.ppm are the same.
13+
Running two_spheres
14+
image width = 400 height = 225
15+
block size = (16, 16) grid size = (25, 15)
16+
Start rendering by GPU.
17+
Done.
18+
two_spheres_gpu.ppm and two_spheres_ref.ppm are the same.
19+
Running two_perlin_spheres
20+
image width = 400 height = 225
21+
block size = (16, 16) grid size = (25, 15)
22+
Start rendering by GPU.
23+
Done.
24+
two_perlin_spheres_gpu.ppm and two_perlin_spheres_ref.ppm are the same.
25+
Running simple_light
26+
image width = 400 height = 225
27+
block size = (16, 16) grid size = (25, 15)
28+
Start rendering by GPU.
29+
Done.
30+
simple_light_gpu.ppm and simple_light_ref.ppm are the same.
31+
Running random_spheres
32+
image width = 400 height = 225
33+
block size = (16, 16) grid size = (25, 15)
34+
Start rendering by GPU.
35+
Done.
36+
random_spheres_gpu.ppm and random_spheres_ref.ppm are the same.
37+
Running cornell_box
38+
image width = 600 height = 600
39+
block size = (16, 16) grid size = (38, 38)
40+
Start rendering by GPU.
41+
Done.
42+
cornell_box_gpu.ppm and cornell_box_ref.ppm are the same.
43+
Running cornell_smoke
44+
image width = 600 height = 600
45+
block size = (16, 16) grid size = (38, 38)
46+
Start rendering by GPU.
47+
Done.
48+
cornell_smoke_gpu.ppm and cornell_smoke_ref.ppm are the same.
49+
exit 0
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,67 @@
1+
#pragma once
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
#include "hipUtils.h"
8+
#include <vector>
9+
10+
template <typename T> class DeviceArray {
11+
public:
12+
// Default constructor
13+
DeviceArray() : DeviceData(nullptr) {}
14+
15+
// Constructor to initialize both host and device data
16+
DeviceArray(size_t N) : HostData(N), DeviceData(nullptr) {
17+
checkHIP(hipMalloc((void **)&DeviceData, N * sizeof(T)),
18+
"Unable to allocate device memory");
19+
}
20+
21+
// Destructor to free device memory
22+
~DeviceArray() { (void)hipFree(DeviceData); }
23+
24+
// Copy data from host to device
25+
void toDevice() {
26+
checkHIP(hipMemcpy(DeviceData, HostData.data(), HostData.size() * sizeof(T),
27+
hipMemcpyHostToDevice),
28+
"Unable to copy data from host to device");
29+
}
30+
31+
// Copy data from device to host
32+
void toHost() {
33+
checkHIP(hipMemcpy(HostData.data(), DeviceData, HostData.size() * sizeof(T),
34+
hipMemcpyDeviceToHost),
35+
"Unable to copy data from device to host");
36+
}
37+
38+
// Resize the array
39+
void resize(size_t N) {
40+
size_t oldSize = HostData.size();
41+
HostData.resize(N);
42+
43+
T *newDeviceData;
44+
checkHIP(hipMalloc((void **)&newDeviceData, N * sizeof(T)),
45+
"Unable to allocate new device memory during resize");
46+
47+
if (DeviceData && oldSize > 0) {
48+
size_t copySize = std::min(oldSize, N) * sizeof(T);
49+
checkHIP(hipMemcpy(newDeviceData, DeviceData, copySize,
50+
hipMemcpyDeviceToDevice),
51+
"Unable to copy data within device during resize");
52+
hipFree(DeviceData);
53+
}
54+
55+
DeviceData = newDeviceData;
56+
}
57+
58+
// Get pointer to host data
59+
T *getHostPtr() { return HostData.data(); }
60+
61+
// Get pointer to device data
62+
T *getDevicePtr() { return DeviceData; }
63+
64+
private:
65+
std::vector<T> HostData; // Host data
66+
T *DeviceData; // Pointer to device data
67+
};
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,184 @@
1+
#pragma once
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
#include <algorithm>
8+
#include <fstream>
9+
#include <iostream>
10+
#include <vector>
11+
12+
#include "color.h"
13+
// Assuming 'color' is a class or struct already defined
14+
// with overloaded operator+ and a method to output the color data
15+
16+
class PPMImageFile {
17+
private:
18+
std::string filename;
19+
int image_width, image_height;
20+
std::vector<color> data;
21+
bool IsNormalized;
22+
23+
public:
24+
PPMImageFile(const std::string &file_name, int width = 0, int height = 0)
25+
: filename(file_name), image_width(width), image_height(height),
26+
IsNormalized(false) {
27+
data.resize(width * height);
28+
}
29+
30+
color *getHostPtr() { return data.data(); }
31+
32+
void setData(color *C) {
33+
for (int i = 0, e = image_width * image_height; i != e; ++i)
34+
data[i] = C[i];
35+
}
36+
37+
void normalize() {
38+
for (auto &pixel_color : data) {
39+
auto r = pixel_color.x();
40+
auto g = pixel_color.y();
41+
auto b = pixel_color.z();
42+
43+
// Apply a linear to gamma transform for gamma 2
44+
r = linear_to_gamma(r);
45+
g = linear_to_gamma(g);
46+
b = linear_to_gamma(b);
47+
48+
// Write the translated [0,255] value of each color component.
49+
static const interval intensity(0.000, 0.999);
50+
pixel_color = color(static_cast<int>(256 * intensity.clamp(r)),
51+
static_cast<int>(256 * intensity.clamp(g)),
52+
static_cast<int>(256 * intensity.clamp(b)));
53+
}
54+
IsNormalized = true;
55+
}
56+
57+
bool save() const {
58+
if (!IsNormalized) {
59+
std::cerr
60+
<< "Error: Image is not normalized. Saving the unnormalized image."
61+
<< std::endl;
62+
return false;
63+
}
64+
65+
std::ofstream file(filename);
66+
67+
if (!file) {
68+
std::cerr << "File could not be opened for writing." << std::endl;
69+
return false;
70+
}
71+
72+
// PPM header
73+
file << "P3\n" << image_width << ' ' << image_height << "\n255\n";
74+
75+
// Write each pixel to the file
76+
for (const auto &pixel_color : data) {
77+
file << pixel_color.x() << ' ' << pixel_color.y() << ' '
78+
<< pixel_color.z() << '\n';
79+
}
80+
81+
file.close();
82+
return true;
83+
}
84+
85+
bool load() {
86+
std::ifstream file(filename);
87+
88+
if (!file) {
89+
std::cerr << "File could not be opened for reading." << std::endl;
90+
return false;
91+
}
92+
93+
std::string header;
94+
int max_val;
95+
96+
// Read the header and check format
97+
file >> header;
98+
if (header != "P3") {
99+
std::cerr << "Unsupported file format." << std::endl;
100+
return false;
101+
}
102+
103+
// Read image dimensions and maximum value
104+
file >> image_width >> image_height >> max_val;
105+
106+
// Resize the data vector to hold the image data
107+
data.resize(image_width * image_height);
108+
109+
// Read pixel data
110+
for (auto &pixel_color : data) {
111+
int r, g, b;
112+
file >> r >> g >> b;
113+
pixel_color = color(r, g, b);
114+
}
115+
116+
IsNormalized = true; // Assuming the loaded image is already normalized
117+
file.close();
118+
return true;
119+
}
120+
bool compare(const PPMImageFile &img, double threshold = 1e-3) const {
121+
if (IsNormalized != img.IsNormalized) {
122+
std::cerr << "Cannot compare " << filename << " and " << img.filename
123+
<< " because one is normalized and the other is not."
124+
<< std::endl;
125+
return false;
126+
}
127+
128+
if (image_width != img.image_width || image_height != img.image_height) {
129+
std::cerr << "Images dimensions do not match." << std::endl;
130+
return false;
131+
}
132+
if (IsNormalized) {
133+
threshold *= 255.0;
134+
}
135+
136+
struct Difference {
137+
double value;
138+
int x, y;
139+
};
140+
141+
std::vector<Difference> topDifferences;
142+
bool anySignificantDifference = false;
143+
144+
for (int y = 0; y < image_height; ++y) {
145+
for (int x = 0; x < image_width; ++x) {
146+
const color &c1 = data[y * image_width + x];
147+
const color &c2 = img.data[y * image_width + x];
148+
149+
double diff =
150+
std::max({std::abs(c1.x() - c2.x()), std::abs(c1.y() - c2.y()),
151+
std::abs(c1.z() - c2.z())});
152+
153+
if (diff > threshold) {
154+
anySignificantDifference = true;
155+
156+
if (topDifferences.size() < 10) {
157+
topDifferences.push_back({diff, x, y});
158+
std::sort(topDifferences.begin(), topDifferences.end(),
159+
[](const Difference &a, const Difference &b) {
160+
return a.value > b.value;
161+
});
162+
}
163+
}
164+
}
165+
}
166+
167+
if (anySignificantDifference) {
168+
// Output top differences
169+
std::cout << "Top Differences between " << filename << " and "
170+
<< img.filename << ":\n";
171+
for (const auto &diff : topDifferences) {
172+
std::cout << "Location (" << diff.x << ", " << diff.y << "), "
173+
<< "Difference: " << diff.value << ", " << filename << ": "
174+
<< data[diff.y * image_width + diff.x].toString() << ", "
175+
<< img.filename << ": "
176+
<< img.data[diff.y * image_width + diff.x].toString() << "\n";
177+
}
178+
return false;
179+
} else {
180+
std::cout << filename << " and " << img.filename << " are the same.\n";
181+
return true;
182+
}
183+
}
184+
};

0 commit comments

Comments
 (0)