Skip to content

Commit

Permalink
Merge remote-tracking branch 'upstream/main'
Browse files Browse the repository at this point in the history
  • Loading branch information
mysterymath committed Apr 3, 2024
2 parents f20ae5a + 95a9f0d commit 04128f2
Show file tree
Hide file tree
Showing 101 changed files with 3,113,479 additions and 1,836 deletions.
24 changes: 21 additions & 3 deletions External/HIP/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -12,8 +12,26 @@ macro(create_local_hip_tests VariantSuffix)
list(APPEND HIP_LOCAL_TESTS empty)
list(APPEND HIP_LOCAL_TESTS with-fopenmp)
list(APPEND HIP_LOCAL_TESTS saxpy)
list(APPEND HIP_LOCAL_TESTS InOneWeekend)
list(APPEND HIP_LOCAL_TESTS TheNextWeek)

# Copy files needed for ray-tracing tests.
file(GLOB IMAGE_FILES "workload/ray-tracing/images/*.jpg" "workload/ray-tracing/images/*.png")
file(COPY ${IMAGE_FILES} DESTINATION "${CMAKE_CURRENT_BINARY_DIR}")

foreach(_hip_test IN LISTS HIP_LOCAL_TESTS)
create_one_local_test(${_hip_test} ${_hip_test}.hip
set(test_source "${_hip_test}.hip")

if(_hip_test STREQUAL "TheNextWeek" OR _hip_test STREQUAL "InOneWeekend")
file(GLOB REF_PPM_FILES "workload/ray-tracing/${_hip_test}/*.ppm")
file(COPY ${REF_PPM_FILES} DESTINATION "${CMAKE_CURRENT_BINARY_DIR}")
set(test_source "workload/ray-tracing/${_hip_test}/main.cc")
# need -mfma to enable FMA in host code
set_source_files_properties(${test_source} PROPERTIES
COMPILE_FLAGS "-xhip -mfma")
endif()

create_one_local_test(${_hip_test} ${test_source}
${VariantOffload} ${VariantSuffix}
"${VariantCPPFLAGS}" "${VariantLibs}")
endforeach()
Expand Down Expand Up @@ -64,15 +82,15 @@ macro(create_hip_tests)
COMMENT "Run all simple HIP tests")

if(NOT AMDGPU_ARCHS)
list(APPEND AMDGPU_ARCHS "gfx906;gfx90a;gfx1030;gfx1100")
list(APPEND AMDGPU_ARCHS "gfx906;gfx90a;gfx1030;gfx1100;native")
endif()

foreach(_RocmPath ${ROCM_PATHS})
get_version(_RocmVersion ${_RocmPath})
set(_HIP_Suffix "hip-${_RocmVersion}")
# Set up HIP test flags
set(_HIP_CPPFLAGS --rocm-path=${_RocmPath})
set(_HIP_LDFLAGS --rocm-path=${_RocmPath} --hip-link -rtlib=compiler-rt -frtlib-add-rpath)
set(_HIP_LDFLAGS --rocm-path=${_RocmPath} --hip-link -rtlib=compiler-rt -unwindlib=libgcc -frtlib-add-rpath)

# Unset these for each iteration of rocm path.
set(_ArchFlags)
Expand Down
6 changes: 6 additions & 0 deletions External/HIP/InOneWeekend.reference_output
Original file line number Diff line number Diff line change
@@ -0,0 +1,6 @@
image width = 1200 height = 675
block size = (16, 16) grid size = (75, 43)
Start rendering by GPU.
Done.
gpu.ppm and ref.ppm are the same.
exit 0
49 changes: 49 additions & 0 deletions External/HIP/TheNextWeek.reference_output
Original file line number Diff line number Diff line change
@@ -0,0 +1,49 @@
Running quads
image width = 400 height = 400
block size = (16, 16) grid size = (25, 25)
Start rendering by GPU.
Done.
quads_gpu.ppm and quads_ref.ppm are the same.
Running earth
image width = 400 height = 225
block size = (16, 16) grid size = (25, 15)
Start rendering by GPU.
Done.
earth_gpu.ppm and earth_ref.ppm are the same.
Running two_spheres
image width = 400 height = 225
block size = (16, 16) grid size = (25, 15)
Start rendering by GPU.
Done.
two_spheres_gpu.ppm and two_spheres_ref.ppm are the same.
Running two_perlin_spheres
image width = 400 height = 225
block size = (16, 16) grid size = (25, 15)
Start rendering by GPU.
Done.
two_perlin_spheres_gpu.ppm and two_perlin_spheres_ref.ppm are the same.
Running simple_light
image width = 400 height = 225
block size = (16, 16) grid size = (25, 15)
Start rendering by GPU.
Done.
simple_light_gpu.ppm and simple_light_ref.ppm are the same.
Running random_spheres
image width = 400 height = 225
block size = (16, 16) grid size = (25, 15)
Start rendering by GPU.
Done.
random_spheres_gpu.ppm and random_spheres_ref.ppm are the same.
Running cornell_box
image width = 600 height = 600
block size = (16, 16) grid size = (38, 38)
Start rendering by GPU.
Done.
cornell_box_gpu.ppm and cornell_box_ref.ppm are the same.
Running cornell_smoke
image width = 600 height = 600
block size = (16, 16) grid size = (38, 38)
Start rendering by GPU.
Done.
cornell_smoke_gpu.ppm and cornell_smoke_ref.ppm are the same.
exit 0
67 changes: 67 additions & 0 deletions External/HIP/workload/ray-tracing/InOneWeekend/DeviceArray.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,67 @@
#pragma once
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
#include "hipUtils.h"
#include <vector>

template <typename T> class DeviceArray {
public:
// Default constructor
DeviceArray() : DeviceData(nullptr) {}

// Constructor to initialize both host and device data
DeviceArray(size_t N) : HostData(N), DeviceData(nullptr) {
checkHIP(hipMalloc((void **)&DeviceData, N * sizeof(T)),
"Unable to allocate device memory");
}

// Destructor to free device memory
~DeviceArray() { (void)hipFree(DeviceData); }

// Copy data from host to device
void toDevice() {
checkHIP(hipMemcpy(DeviceData, HostData.data(), HostData.size() * sizeof(T),
hipMemcpyHostToDevice),
"Unable to copy data from host to device");
}

// Copy data from device to host
void toHost() {
checkHIP(hipMemcpy(HostData.data(), DeviceData, HostData.size() * sizeof(T),
hipMemcpyDeviceToHost),
"Unable to copy data from device to host");
}

// Resize the array
void resize(size_t N) {
size_t oldSize = HostData.size();
HostData.resize(N);

T *newDeviceData;
checkHIP(hipMalloc((void **)&newDeviceData, N * sizeof(T)),
"Unable to allocate new device memory during resize");

if (DeviceData && oldSize > 0) {
size_t copySize = std::min(oldSize, N) * sizeof(T);
checkHIP(hipMemcpy(newDeviceData, DeviceData, copySize,
hipMemcpyDeviceToDevice),
"Unable to copy data within device during resize");
hipFree(DeviceData);
}

DeviceData = newDeviceData;
}

// Get pointer to host data
T *getHostPtr() { return HostData.data(); }

// Get pointer to device data
T *getDevicePtr() { return DeviceData; }

private:
std::vector<T> HostData; // Host data
T *DeviceData; // Pointer to device data
};
184 changes: 184 additions & 0 deletions External/HIP/workload/ray-tracing/InOneWeekend/PPMImageFile.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,184 @@
#pragma once
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
#include <algorithm>
#include <fstream>
#include <iostream>
#include <vector>

#include "color.h"
// Assuming 'color' is a class or struct already defined
// with overloaded operator+ and a method to output the color data

class PPMImageFile {
private:
std::string filename;
int image_width, image_height;
std::vector<color> data;
bool IsNormalized;

public:
PPMImageFile(const std::string &file_name, int width = 0, int height = 0)
: filename(file_name), image_width(width), image_height(height),
IsNormalized(false) {
data.resize(width * height);
}

color *getHostPtr() { return data.data(); }

void setData(color *C) {
for (int i = 0, e = image_width * image_height; i != e; ++i)
data[i] = C[i];
}

void normalize() {
for (auto &pixel_color : data) {
auto r = pixel_color.x();
auto g = pixel_color.y();
auto b = pixel_color.z();

// Apply a linear to gamma transform for gamma 2
r = linear_to_gamma(r);
g = linear_to_gamma(g);
b = linear_to_gamma(b);

// Write the translated [0,255] value of each color component.
static const interval intensity(0.000, 0.999);
pixel_color = color(static_cast<int>(256 * intensity.clamp(r)),
static_cast<int>(256 * intensity.clamp(g)),
static_cast<int>(256 * intensity.clamp(b)));
}
IsNormalized = true;
}

bool save() const {
if (!IsNormalized) {
std::cerr
<< "Error: Image is not normalized. Saving the unnormalized image."
<< std::endl;
return false;
}

std::ofstream file(filename);

if (!file) {
std::cerr << "File could not be opened for writing." << std::endl;
return false;
}

// PPM header
file << "P3\n" << image_width << ' ' << image_height << "\n255\n";

// Write each pixel to the file
for (const auto &pixel_color : data) {
file << pixel_color.x() << ' ' << pixel_color.y() << ' '
<< pixel_color.z() << '\n';
}

file.close();
return true;
}

bool load() {
std::ifstream file(filename);

if (!file) {
std::cerr << "File could not be opened for reading." << std::endl;
return false;
}

std::string header;
int max_val;

// Read the header and check format
file >> header;
if (header != "P3") {
std::cerr << "Unsupported file format." << std::endl;
return false;
}

// Read image dimensions and maximum value
file >> image_width >> image_height >> max_val;

// Resize the data vector to hold the image data
data.resize(image_width * image_height);

// Read pixel data
for (auto &pixel_color : data) {
int r, g, b;
file >> r >> g >> b;
pixel_color = color(r, g, b);
}

IsNormalized = true; // Assuming the loaded image is already normalized
file.close();
return true;
}
bool compare(const PPMImageFile &img, double threshold = 1e-3) const {
if (IsNormalized != img.IsNormalized) {
std::cerr << "Cannot compare " << filename << " and " << img.filename
<< " because one is normalized and the other is not."
<< std::endl;
return false;
}

if (image_width != img.image_width || image_height != img.image_height) {
std::cerr << "Images dimensions do not match." << std::endl;
return false;
}
if (IsNormalized) {
threshold *= 255.0;
}

struct Difference {
double value;
int x, y;
};

std::vector<Difference> topDifferences;
bool anySignificantDifference = false;

for (int y = 0; y < image_height; ++y) {
for (int x = 0; x < image_width; ++x) {
const color &c1 = data[y * image_width + x];
const color &c2 = img.data[y * image_width + x];

double diff =
std::max({std::abs(c1.x() - c2.x()), std::abs(c1.y() - c2.y()),
std::abs(c1.z() - c2.z())});

if (diff > threshold) {
anySignificantDifference = true;

if (topDifferences.size() < 10) {
topDifferences.push_back({diff, x, y});
std::sort(topDifferences.begin(), topDifferences.end(),
[](const Difference &a, const Difference &b) {
return a.value > b.value;
});
}
}
}
}

if (anySignificantDifference) {
// Output top differences
std::cout << "Top Differences between " << filename << " and "
<< img.filename << ":\n";
for (const auto &diff : topDifferences) {
std::cout << "Location (" << diff.x << ", " << diff.y << "), "
<< "Difference: " << diff.value << ", " << filename << ": "
<< data[diff.y * image_width + diff.x].toString() << ", "
<< img.filename << ": "
<< img.data[diff.y * image_width + diff.x].toString() << "\n";
}
return false;
} else {
std::cout << filename << " and " << img.filename << " are the same.\n";
return true;
}
}
};
Loading

0 comments on commit 04128f2

Please sign in to comment.