Add HIP test for ray-tracing workload

External/HIP/CMakeLists.txt

@@ -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()
@@ -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)

External/HIP/InOneWeekend.reference_output

@@ -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

External/HIP/TheNextWeek.reference_output

@@ -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

External/HIP/workload/ray-tracing/InOneWeekend/DeviceArray.h

@@ -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
+};

External/HIP/workload/ray-tracing/InOneWeekend/PPMImageFile.h

@@ -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;
+    }
+  }
+};