EXSWCPHIPT-42 - Add HIP RTC support to the test framework (#2719)

* EXSWCPHIPT-42 - Add HIP RTC support to the test framework
* Removed ifdef from hipTestContext class

catch/include/hip_test_common.hh

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2021 - 2021 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2021 - 2022 Advanced Micro Devices, Inc. All rights reserved.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -22,6 +22,7 @@ THE SOFTWARE.
 
 #pragma once
 #include "hip_test_context.hh"
+#include <hip_test_rtc.hh>
 #include <catch.hpp>
 #include <stdlib.h>
 
@@ -171,8 +172,60 @@ static inline void HIP_SKIP_TEST(char const* const reason) noexcept {
   // ctest is setup to parse for "HIP_SKIP_THIS_TEST", at which point it will skip the test.
   std::cout << "Skipping test. Reason: " << reason << '\n' << "HIP_SKIP_THIS_TEST" << std::endl;
 }
+
+/**
+ * @brief Helper template that returns the expected arguments of a kernel.
+ *
+ * @return constexpr std::tuple<FArgs...> the expected arguments of the kernel.
+ */
+template <typename... FArgs> std::tuple<FArgs...> getExpectedArgs(void(FArgs...)){};
+
+/**
+ * @brief Asserts that the types of the arguments of a function match exactly with the types in the
+ * function signature.
+ * This is necessary because HIP RTC does not do implicit casting of the kernel
+ * parameters.
+ * In order to get the kernel function signature, this function should only called when
+ * RTC is disabled.
+ *
+ * @tparam F the kernel function
+ * @tparam Args the parameters that will be passed to the kernel.
+ */
+template <typename F, typename... Args> void validateArguments(F f, Args...) {
+  using expectedArgsTuple = decltype(getExpectedArgs(f));
+  static_assert(std::is_same<expectedArgsTuple, std::tuple<Args...>>::value,
+                "Kernel arguments types must match exactly!");
 }
 
+/**
+ * @brief Launch a kernel using either HIP or HIP RTC.
+ *
+ * @tparam Typenames A list of typenames used by the kernel (unused if the kernel is not a
+ * template).
+ * @tparam K The kernel type. Expects a function or template when RTC is disabled. Expects a
+ * function pointer instead when RTC is enabled.
+ * @tparam Dim Can be either dim3 or int.
+ * @tparam Args A list of kernel arguments to be forwarded.
+ * @param kernel The kernel to be launched (defined in kernels.hh)
+ * @param numBlocks
+ * @param numThreads
+ * @param memPerBlock
+ * @param stream
+ * @param packedArgs A list of kernel arguments to be forwarded.
+ */
+template <typename... Typenames, typename K, typename Dim, typename... Args>
+void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerBlock,
+                  hipStream_t stream, Args&&... packedArgs) {
+#ifndef RTC_TESTING
+    validateArguments(kernel, packedArgs...);
+    kernel<<<numBlocks, numThreads, memPerBlock, stream>>>(std::forward<Args>(packedArgs)...);
+#else
+  launchRTCKernel<Typenames...>(kernel, numBlocks, numThreads, memPerBlock, stream,
+                                std::forward<Args>(packedArgs)...);
+#endif
+}
+}  // namespace HipTest
+
 
 // This must be called in the beginning of image test app's main() to indicate whether image
 // is supported.

catch/include/hip_test_context.hh

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -22,9 +22,11 @@ THE SOFTWARE.
 
 #pragma once
 #include <hip/hip_runtime.h>
+#include <hip/hiprtc.h>
 #include <vector>
 #include <string>
 #include <set>
+#include <unordered_map>
 
 #if defined(_WIN32)
 #define HT_WIN 1
@@ -57,9 +59,9 @@ static int _log_enable = (std::getenv("HT_LOG_ENABLE") ? 1 : 0);
   }
 
 typedef struct Config_ {
-  std::string json_file;              // Json file
-  std::string platform;               // amd/nvidia
-  std::string os;                     // windows/linux
+  std::string json_file;  // Json file
+  std::string platform;   // amd/nvidia
+  std::string os;         // windows/linux
 } Config;
 
 class TestContext {
@@ -69,14 +71,20 @@ class TestContext {
   std::string current_test;
   std::set<std::string> skip_test;
   std::string json_file_;
-  std::vector<std::string>  platform_list_ = {"amd" , "nvidia"};
-  std::vector<std::string>  os_list_ = {"windows", "linux", "all"};
-  std::vector<std::string>  amd_arch_list_ = {};
+  std::vector<std::string> platform_list_ = {"amd", "nvidia"};
+  std::vector<std::string> os_list_ = {"windows", "linux", "all"};
+  std::vector<std::string> amd_arch_list_ = {};
+
+  struct rtcState {
+    hipModule_t module;
+    hipFunction_t kernelFunction;
+  };
+
+  std::unordered_map<std::string, rtcState> compiledKernels{};
 
   Config config_;
   std::string& getJsonFile();
-  std::string substringFound( std::vector<std::string> list,
-                              std::string filename);
+  std::string substringFound(std::vector<std::string> list, std::string filename);
   void detectOS();
   void detectPlatform();
   void fillConfig();
@@ -86,10 +94,11 @@ class TestContext {
   std::string getMatchingConfigFile(std::string config_dir);
   const Config& getConfig() const { return config_; }
 
+
   TestContext(int argc, char** argv);
 
  public:
-  static const TestContext& get(int argc = 0, char** argv = nullptr) {
+  static TestContext& get(int argc = 0, char** argv = nullptr) {
     static TestContext instance(argc, argv);
     return instance;
   }
@@ -103,6 +112,34 @@ class TestContext {
   const std::string& getCurrentTest() const { return current_test; }
   std::string currentPath() const;
 
+  /**
+   * @brief Unload all loaded modules.
+   * Note: This function needs to be called at the end of each test that uses RTC.
+   *       It is not possible to unload the loaded modules without adding explicit code to the end
+   * of each test. This function exists only to provide a clean way to exit a test when using RTC.
+   *       However, not unloading a module explicitly shouldn't have any effect on the outcome of
+   * the test.
+   */
+  void cleanContext();
+
+  /**
+   * @brief Keeps track of all the already compiled rtc kernels.
+   *
+   * @param kernelNameExpression The name expression (e.g. hipTest::vectorADD<float>).
+   * @param loadedModule  The loaded module.
+   * @param kernelFunction The hipFunction that will be used to run the kernel in the future.
+   */
+  void trackRtcState(std::string kernelNameExpression, hipModule_t loadedModule,
+                     hipFunction_t kernelFunction);
+
+  /**
+   * @brief Get the already compiled hip rtc kernel function if it exists.
+   *
+   * @param kernelNameExpression The name expression (e.g. hipTest::vectorADD<float>).
+   * @return the hipFunction if it exists. nullptr otherwise
+   */
+  hipFunction_t getFunction(const std::string kernelNameExpression);
+
   TestContext(const TestContext&) = delete;
   void operator=(const TestContext&) = delete;
 };

catch/include/hip_test_rtc.hh

@@ -0,0 +1,279 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#pragma once
+#include <hip/hip_runtime.h>
+#include <hip/hiprtc.h>
+#include <kernel_mapping.hh>
+#include <catch.hpp>
+#include <string>
+#include <vector>
+#include <iostream>
+#include <fstream>
+#include <sstream>
+#include <set>
+#include "hip/hip_runtime_api.h"
+#include "hip_test_context.hh"
+
+namespace HipTest {
+
+struct KernelArgument {
+  const void* ptr;
+  size_t sizeRequirement;
+  size_t alignmentRequirement;
+};
+
+/**
+ * @brief Reconstructs the name expression for the kernel.
+ *
+ * @param kernelName the name of the kernel (e.g. "HipTest::VectorADD")
+ * @param typenames the typenames used by this kernel (e.g. "float").
+ * @return std::string the reconstructed expression (e.g. "VectorADD<float>""). Returns kernelName
+ * instead if the kernel is not a template.
+ */
+inline std::string reconstructExpression(std::string& kernelName,
+                                         std::vector<std::string>& typenames) {
+  std::string kernelExpression = kernelName;
+  if (typenames.size() > 0) {
+    kernelExpression += "<" + typenames[0];
+    for (size_t i = 1; i < typenames.size(); ++i) {
+      kernelExpression += "," + typenames[i];
+    }
+    kernelExpression += ">";
+  }
+
+  return kernelExpression;
+}
+
+/**
+ * @brief Packs the kernel arguments into the format expected by hipModuleLaunchKernel
+ *
+ * @param args list of arguments for the kernel and their alignemnt requirements.
+ * @return std::vector<char> the packed arguments ready to be passed on to hipModuleLaunchKernel
+ */
+inline std::vector<char> alignArguments(std::vector<KernelArgument>& args) {
+  std::vector<char> alignedArguments{};
+  int count = 0;
+  for (auto& arg : args) {
+    const char* argPtr{reinterpret_cast<const char*>(arg.ptr)};
+
+    /*
+     * Details about the padding formula can be found at:
+     * https://en.wikipedia.org/wiki/Data_structure_alignment#Data_structure_padding
+     */
+    int paddingNeeded = -count & (arg.alignmentRequirement - 1);
+    alignedArguments.insert(std::end(alignedArguments), paddingNeeded, 0);
+    count += paddingNeeded;
+
+    alignedArguments.insert(std::end(alignedArguments), argPtr, argPtr + arg.sizeRequirement);
+    count += arg.sizeRequirement;
+  }
+  return alignedArguments;
+}
+
+inline std::vector<char> getKernelCode(hiprtcProgram& rtcProgram) {
+  size_t codeSize;
+  REQUIRE(HIPRTC_SUCCESS == hiprtcGetCodeSize(rtcProgram, &codeSize));
+
+  std::vector<char> code(codeSize);
+  REQUIRE(HIPRTC_SUCCESS == hiprtcGetCode(rtcProgram, code.data()));
+
+  return code;
+}
+
+/**
+ * @brief Compiles a kernel using HIP RTC
+ *
+ * @param rtcKernel the name of the kernel to compile.
+ * @param kernelNameExpression the name expression to be added to the RTC program (e.g.
+ * HipTest::VectorADD<float>)
+ * @return hiprtcProgram the compiled rtc program.
+ */
+inline hiprtcProgram compileRTC(std::string& rtcKernel, std::string& kernelNameExpression) {
+  std::string fileName = mapKernelToFileName.at(rtcKernel);
+  std::string filePath{KERNELS_PATH + fileName};
+
+  INFO("Opening Kernel File: " << filePath);
+  std::ifstream kernelFile{filePath};
+  REQUIRE(kernelFile.is_open());
+
+  std::stringstream stringStream;
+  std::string line;
+  while (getline(kernelFile, line)) {
+    /* Skip the include directive since it is not part of the kernel */
+    if (line.find("#include") != std::string::npos) {
+      continue;
+    }
+    stringStream << line << '\n';
+  }
+  kernelFile.close();
+
+  std::string kernelCode{stringStream.str()};
+  INFO("RTC Kernel Code:\n" << kernelCode)
+
+  hiprtcProgram rtcProgram;
+  hiprtcCreateProgram(&rtcProgram, kernelCode.c_str(), (fileName + ".cu").c_str(), 0, nullptr,
+                      nullptr);
+
+  std::vector<const char*> options{};
+#ifdef __HIP_PLATFORM_AMD__
+
+  int deviceCount;
+  REQUIRE(hipSuccess == hipGetDeviceCount(&deviceCount));
+
+  std::set<std::string> architectures{};
+  for (int i = 0; i < deviceCount; ++i) {
+    hipDeviceProp_t props;
+    REQUIRE(hipSuccess == hipGetDeviceProperties(&props, i));
+    architectures.insert(std::string{"--gpu-architecture="} + props.gcnArchName);
+  }
+
+  for (auto& architecture : architectures) {
+    options.push_back(architecture.c_str());
+  }
+#else
+  options.push_back("--fmad=false");
+#endif
+
+  REQUIRE(HIPRTC_SUCCESS == hiprtcAddNameExpression(rtcProgram, kernelNameExpression.c_str()));
+  REQUIRE(HIPRTC_SUCCESS == hiprtcCompileProgram(rtcProgram, 1, options.data()));
+
+  return rtcProgram;
+}
+
+/**
+ * @brief Get a typename as a string
+ *
+ * @tparam T The typename
+ * @return std::string the string representation of T
+ */
+template <typename T> std::string getTypeName() {
+  std::string name, prefix, suffix;
+
+
+#ifdef __clang__
+  name = __PRETTY_FUNCTION__;
+  prefix = "std::string HipTest::getTypeName() [T = ";
+  suffix = "]";
+#elif defined(__GNUC__)
+  name = __PRETTY_FUNCTION__;
+  prefix = "std::string HipTest::getTypeName() [with T = ";
+  suffix = "; std::string = std::__cxx11::basic_string<char>]";
+#elif defined(_MSC_VER)
+  name = __FUNCSIG__;
+  prefix = "std::string __cdecl HipTest::getTypeName<";
+  suffix = ">(void)";
+#endif
+
+  return name.substr(prefix.size(), name.rfind(suffix) - prefix.size());
+}
+
+/**
+ * @brief Tells the user that the kernels are using HIP RTC. Prints only once per test.
+ *
+ */
+static inline void printInfo() {
+  static bool alreadyPrinted{false};
+
+  if (!alreadyPrinted) {
+    std::cout << "INFO: This test is running using HIP RTC to compile and run the kernels."
+              << std::endl;
+    alreadyPrinted = true;
+  }
+}
+
+/**
+ * @brief Compiles and launches a kernel using HIP RTC
+ *
+ * @tparam Typenames A list of typenames used by the kernel (unused if the kernel is not a
+ * template).
+ * @tparam Args A list of kernel arguments to be forwarded.
+ * @param getKernelName A function wrapper that returns the name of the kernel to launch (check
+ * kernels.hh for more info)
+ * @param numBlocks
+ * @param numThreads
+ * @param memPerBlock
+ * @param stream
+ * @param packedArgs A list of kernel arguments to be forwarded.
+ */
+template <typename... Typenames, typename... Args>
+void launchRTCKernel(std::string (*getKernelName)(), dim3 numBlocks, dim3 numThreads,
+                     std::uint32_t memPerBlock, hipStream_t stream, Args&&... packedArgs) {
+
+  printInfo();
+  TestContext& testContext = TestContext::get();
+  std::string kernelName = (*getKernelName)();
+
+  std::vector<std::string> kernelTypenames{std::string(HipTest::getTypeName<Typenames>())...};
+  std::string kernelExpression = reconstructExpression(kernelName, kernelTypenames);
+
+  if (testContext.getFunction(kernelExpression) == nullptr) {
+    hiprtcProgram rtcProgram{compileRTC(kernelName, kernelExpression)};
+    std::vector<char> compiledCode{getKernelCode(rtcProgram)};
+
+    hipModule_t module;
+
+    REQUIRE(hipSuccess == hipModuleLoadData(&module, compiledCode.data()));
+
+      hipFunction_t kernelFunction;
+
+    const char* loweredName;
+    REQUIRE(HIPRTC_SUCCESS ==
+          hiprtcGetLoweredName(rtcProgram, kernelExpression.c_str(), &loweredName));
+    REQUIRE(hipSuccess == hipModuleGetFunction(&kernelFunction, module, loweredName));
+
+    /* After obtaining the kernelFunction, the program is no longer needed. So it can be destroyed */
+    REQUIRE(HIPRTC_SUCCESS == hiprtcDestroyProgram(&rtcProgram));
+
+    testContext.trackRtcState(kernelExpression, module, kernelFunction);
+  }
+
+  hipFunction_t kernelFunction = testContext.getFunction(kernelExpression);
+
+  std::vector<KernelArgument> args = {
+      {reinterpret_cast<const void*>(&packedArgs), sizeof(Args), alignof(Args)}...};
+
+  std::vector<char> alignedArguments{alignArguments(args)};
+  size_t argumentsSize{alignedArguments.size()};
+
+  void* config_array[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, alignedArguments.data(),
+                          HIP_LAUNCH_PARAM_BUFFER_SIZE, reinterpret_cast<void*>(&argumentsSize),
+                          HIP_LAUNCH_PARAM_END};
+
+  REQUIRE(hipSuccess ==
+          hipModuleLaunchKernel(kernelFunction, numBlocks.x, numBlocks.y, numBlocks.z, numThreads.x,
+                                numThreads.y, numThreads.z, memPerBlock, stream, nullptr,
+                                config_array));
+}
+
+/**
+ * @brief Template overload for when numBlocks and numThreads is an integer.
+ *
+ */
+template <typename... Typenames, typename... Args>
+void launchRTCKernel(std::string kernelName, int numBlocks, int numThreads,
+                     std::uint32_t memPerBlock, hipStream_t stream, Args&&... packedArgs) {
+  launchRTCKernel<Typenames...>(kernelName, dim3(numBlocks), dim3(numThreads), memPerBlock, stream,
+                                std::forward<Args>(packedArgs)...);
+}
+
+}  // namespace HipTest

catch/include/kernel_mapping.hh

@@ -0,0 +1,27 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#pragma once
+
+#include <map>
+
+const std::map<std::string, std::string> mapKernelToFileName{
+  {"Set", "Set.cpp"},
+  {"HipTest::vectorADD", "vectorADD.inl"},
+};

catch/include/kernels.hh

@@ -0,0 +1,55 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#pragma once
+
+#include <hip_test_common.hh>
+#include <map>
+
+#ifndef RTC_TESTING
+
+__global__ void Set(int* Ad, int val);
+
+/* Kernel Templates */
+#include "vectorADD.inl"
+
+#else
+
+/*
+ * Wrapper Macros that create a string representation of the kernel name.
+ * In the case of kernel templates, a variadic template is used to ensure compatibility with
+ * the launchKernel template when RTC is not enabled. If the kernel is inside a namespace, use the
+ * "_NS" version of the Macro.
+ */
+#define FUNCTION_WRAPPER(param)                                                                    \
+  std::string param() { return #param; }
+#define TEMPLATE_WRAPPER(param)                                                                    \
+  template <typename...> std::string param() { return #param; }
+#define FUNCTION_WRAPPER_NS(param, namespace)                                                      \
+  std::string param() { return #namespace "::" #param; }
+#define TEMPLATE_WRAPPER_NS(param, namespace)                                                      \
+  template <typename...> std::string param() { return #namespace "::" #param; }
+
+FUNCTION_WRAPPER(Set);
+
+namespace HipTest {
+TEMPLATE_WRAPPER_NS(vectorADD, HipTest);
+}
+
+#endif