diff --git a/tests/README.md b/tests/README.md index 33912a5726..1b1e54a2e3 100644 --- a/tests/README.md +++ b/tests/README.md @@ -147,6 +147,22 @@ For example, Here "-C performance" indicate the "performance" configuration of ctest. ``` +### RTC Testing + +To enable RTC testing, cmake needs to be passed the DRTC_TESTING=1 options. + +When this option is passed, all tests that support this functionality will be run using HIP RTC to compile and run. + +To enable HIP RTC support for a specific test: + 1 - Move all its kernels to tests/catch/kernels (one file per kernel) + 2 - Update tests/catch/kernels/CMakeLists.txt + 3 - Update tests/catch/include/kernels.hh + 4 - Update tests/catch/include/kernel_mapping.hh + 5 - Include kernels.hh + 6 - Call hipTest::launchKernel() function instead of hipLaunchKernelGGL() + +Note: HIP RTC does not do implicit casting of kernel parameters. This requires the test writer to explicitly do all the casting before running the kernel. The code will not compile otherwise. + ### If a test fails - how to debug a test Find the test and commandline that fail: diff --git a/tests/catch/CMakeLists.txt b/tests/catch/CMakeLists.txt index 40a61fd6f7..1a272d03aa 100644 --- a/tests/catch/CMakeLists.txt +++ b/tests/catch/CMakeLists.txt @@ -104,10 +104,17 @@ set(CATCH2_INCLUDE ${CATCH2_PATH}/cmake/Catch2/catch_include.cmake.in) include_directories( ${CATCH2_PATH} "./include" + "./kernels" ${HIP_PATH}/include ${JSON_PARSER} ) +option(RTC_TESTING "Run tests using HIP RTC to compile the kernels" OFF) +if (RTC_TESTING) + add_definitions(-DRTC_TESTING=ON) +endif() +add_definitions(-DKERNELS_PATH="${CMAKE_CURRENT_SOURCE_DIR}/kernels/") + file(COPY ./hipTestMain/config DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/hipTestMain) file(COPY ./external/Catch2/cmake/Catch2/CatchAddTests.cmake DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/script) set(ADD_SCRIPT_PATH ${CMAKE_CURRENT_BINARY_DIR}/script/CatchAddTests.cmake) @@ -177,6 +184,7 @@ add_custom_target(build_tests) # Tests folder add_subdirectory(unit) add_subdirectory(ABM) +add_subdirectory(kernels) add_subdirectory(hipTestMain) add_subdirectory(stress) add_subdirectory(TypeQualifiers) diff --git a/tests/catch/external/Catch2/cmake/Catch2/Catch.cmake b/tests/catch/external/Catch2/cmake/Catch2/Catch.cmake index 5287241143..62e4984afc 100644 --- a/tests/catch/external/Catch2/cmake/Catch2/Catch.cmake +++ b/tests/catch/external/Catch2/cmake/Catch2/Catch.cmake @@ -203,7 +203,16 @@ function(hip_add_exe_to_target) "${list_args}" ) # Create shared lib of all tests + if(NOT RTC_TESTING) + add_executable(${_NAME} EXCLUDE_FROM_ALL ${_TEST_SRC} $ $) + else () add_executable(${_NAME} EXCLUDE_FROM_ALL ${_TEST_SRC} $) + if(HIP_PLATFORM STREQUAL "amd") + target_link_libraries(${_NAME} hiprtc) + else() + target_link_libraries(${_NAME} nvrtc) + endif() + endif() catch_discover_tests(${_NAME} PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST") if(UNIX) set(_LINKER_LIBS ${_LINKER_LIBS} stdc++fs) diff --git a/tests/catch/hipTestMain/hip_test_context.cc b/tests/catch/hipTestMain/hip_test_context.cc index ec55e2395e..a6e3a08609 100644 --- a/tests/catch/hipTestMain/hip_test_context.cc +++ b/tests/catch/hipTestMain/hip_test_context.cc @@ -23,10 +23,10 @@ void TestContext::detectPlatform() { std::string TestContext::substringFound(std::vector list, std::string filename) { std::string match = ""; - for(unsigned int i = 0; i < list.size() ; i++) { + for (unsigned int i = 0; i < list.size(); i++) { if (filename.find(list.at(i)) != std::string::npos) { - match = list.at(i); - break; + match = list.at(i); + break; } } return match; @@ -35,20 +35,19 @@ std::string TestContext::substringFound(std::vector list, std::stri std::string TestContext::getMatchingConfigFile(std::string config_dir) { std::string configFileToUse; - for(auto& p: fs::recursive_directory_iterator(config_dir)) { + for (auto& p : fs::recursive_directory_iterator(config_dir)) { fs::path filename = p.path(); std::string cur_arch = "TODO"; - std::string arch = substringFound(amd_arch_list_,filename.filename().string()); - std::string platform = substringFound(platform_list_,filename.filename().string()); - std::string os = substringFound(os_list_,filename.filename().string()); + std::string arch = substringFound(amd_arch_list_, filename.filename().string()); + std::string platform = substringFound(platform_list_, filename.filename().string()); + std::string os = substringFound(os_list_, filename.filename().string()); // if arch found then use that exit from loop - if(arch == cur_arch) { + if (arch == cur_arch) { configFileToUse = filename.string(); break; - // match the platform/os and continue to look - } else if((platform == config_.platform) && - (os == config_.os || os == "all")) { - configFileToUse = filename.string(); + // match the platform/os and continue to look + } else if ((platform == config_.platform) && (os == config_.os || os == "all")) { + configFileToUse = filename.string(); } } return configFileToUse; @@ -60,10 +59,10 @@ std::string& TestContext::getJsonFile() { config_dir = config_dir.parent_path(); int levels = 0; bool configFolderFound = false; - std::vector configList; + std::vector configList; std::string configFile; // check a max of 5 levels down the executable path - while(levels < 5) { + while (levels < 5) { fs::path temp_path = config_dir; temp_path /= "hipTestMain"; temp_path /= "config"; @@ -185,7 +184,7 @@ bool TestContext::parseJsonFile() { return false; } - const picojson::object &o = v.get(); + const picojson::object& o = v.get(); for (picojson::object::const_iterator i = o.begin(); i != o.end(); ++i) { // Processing for DisabledTests if (i->first == "DisabledTests") { @@ -196,7 +195,7 @@ bool TestContext::parseJsonFile() { for (auto ai = val.begin(); ai != val.end(); ai++) { std::string tmp = ai->get(); std::string newRegexName; - for(const auto &c : tmp) { + for (const auto& c : tmp) { if (c == '*') newRegexName += ".*"; else @@ -209,3 +208,25 @@ bool TestContext::parseJsonFile() { return true; } + +void TestContext::cleanContext() { + for (auto& pair : compiledKernels) { + REQUIRE(hipSuccess == hipModuleUnload(pair.second.module)); + } +} + +void TestContext::trackRtcState(std::string kernelNameExpression, hipModule_t loadedModule, + hipFunction_t kernelFunction) { + rtcState state{loadedModule, kernelFunction}; + compiledKernels[kernelNameExpression] = state; +} + +hipFunction_t TestContext::getFunction(const std::string kernelNameExpression) { + auto it{compiledKernels.find(kernelNameExpression)}; + + if (it != compiledKernels.end()) { + return it->second.kernelFunction; + } else { + return nullptr; + } +} \ No newline at end of file diff --git a/tests/catch/include/hip_test_common.hh b/tests/catch/include/hip_test_common.hh index 496f202b8c..63b8354421 100644 --- a/tests/catch/include/hip_test_common.hh +++ b/tests/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 #include #include @@ -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 the expected arguments of the kernel. + */ +template std::tuple 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 void validateArguments(F f, Args...) { + using expectedArgsTuple = decltype(getExpectedArgs(f)); + static_assert(std::is_same>::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 +void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerBlock, + hipStream_t stream, Args&&... packedArgs) { +#ifndef RTC_TESTING + validateArguments(kernel, packedArgs...); + kernel<<>>(std::forward(packedArgs)...); +#else + launchRTCKernel(kernel, numBlocks, numThreads, memPerBlock, stream, + std::forward(packedArgs)...); +#endif +} +} // namespace HipTest + // This must be called in the beginning of image test app's main() to indicate whether image // is supported. diff --git a/tests/catch/include/hip_test_context.hh b/tests/catch/include/hip_test_context.hh index e6429454ae..2d630b70f7 100644 --- a/tests/catch/include/hip_test_context.hh +++ b/tests/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 +#include #include #include #include +#include #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 skip_test; std::string json_file_; - std::vector platform_list_ = {"amd" , "nvidia"}; - std::vector os_list_ = {"windows", "linux", "all"}; - std::vector amd_arch_list_ = {}; + std::vector platform_list_ = {"amd", "nvidia"}; + std::vector os_list_ = {"windows", "linux", "all"}; + std::vector amd_arch_list_ = {}; + + struct rtcState { + hipModule_t module; + hipFunction_t kernelFunction; + }; + + std::unordered_map compiledKernels{}; Config config_; std::string& getJsonFile(); - std::string substringFound( std::vector list, - std::string filename); + std::string substringFound(std::vector 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). + * @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). + * @return the hipFunction if it exists. nullptr otherwise + */ + hipFunction_t getFunction(const std::string kernelNameExpression); + TestContext(const TestContext&) = delete; void operator=(const TestContext&) = delete; }; diff --git a/tests/catch/include/hip_test_rtc.hh b/tests/catch/include/hip_test_rtc.hh new file mode 100644 index 0000000000..0860ae51db --- /dev/null +++ b/tests/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 +#include +#include +#include +#include +#include +#include +#include +#include +#include +#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""). Returns kernelName + * instead if the kernel is not a template. + */ +inline std::string reconstructExpression(std::string& kernelName, + std::vector& 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 the packed arguments ready to be passed on to hipModuleLaunchKernel + */ +inline std::vector alignArguments(std::vector& args) { + std::vector alignedArguments{}; + int count = 0; + for (auto& arg : args) { + const char* argPtr{reinterpret_cast(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 getKernelCode(hiprtcProgram& rtcProgram) { + size_t codeSize; + REQUIRE(HIPRTC_SUCCESS == hiprtcGetCodeSize(rtcProgram, &codeSize)); + + std::vector 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) + * @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 options{}; +#ifdef __HIP_PLATFORM_AMD__ + + int deviceCount; + REQUIRE(hipSuccess == hipGetDeviceCount(&deviceCount)); + + std::set 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 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]"; +#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 +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 kernelTypenames{std::string(HipTest::getTypeName())...}; + std::string kernelExpression = reconstructExpression(kernelName, kernelTypenames); + + if (testContext.getFunction(kernelExpression) == nullptr) { + hiprtcProgram rtcProgram{compileRTC(kernelName, kernelExpression)}; + std::vector 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 args = { + {reinterpret_cast(&packedArgs), sizeof(Args), alignof(Args)}...}; + + std::vector 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(&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 +void launchRTCKernel(std::string kernelName, int numBlocks, int numThreads, + std::uint32_t memPerBlock, hipStream_t stream, Args&&... packedArgs) { + launchRTCKernel(kernelName, dim3(numBlocks), dim3(numThreads), memPerBlock, stream, + std::forward(packedArgs)...); +} + +} // namespace HipTest diff --git a/tests/catch/include/kernel_mapping.hh b/tests/catch/include/kernel_mapping.hh new file mode 100644 index 0000000000..32ca93b5e6 --- /dev/null +++ b/tests/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 + +const std::map mapKernelToFileName{ + {"Set", "Set.cpp"}, + {"HipTest::vectorADD", "vectorADD.inl"}, +}; \ No newline at end of file diff --git a/tests/catch/include/kernels.hh b/tests/catch/include/kernels.hh new file mode 100644 index 0000000000..ab2a92bfdb --- /dev/null +++ b/tests/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 +#include + +#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 std::string param() { return #param; } +#define FUNCTION_WRAPPER_NS(param, namespace) \ + std::string param() { return #namespace "::" #param; } +#define TEMPLATE_WRAPPER_NS(param, namespace) \ + template std::string param() { return #namespace "::" #param; } + +FUNCTION_WRAPPER(Set); + +namespace HipTest { +TEMPLATE_WRAPPER_NS(vectorADD, HipTest); +} + +#endif diff --git a/tests/catch/kernels/CMakeLists.txt b/tests/catch/kernels/CMakeLists.txt new file mode 100644 index 0000000000..91e1ab69ff --- /dev/null +++ b/tests/catch/kernels/CMakeLists.txt @@ -0,0 +1,8 @@ +if(NOT RTC_TESTING) + set(TEST_SRC + Set.cpp + ) + + add_library(KERNELS EXCLUDE_FROM_ALL OBJECT ${TEST_SRC}) + target_compile_options(KERNELS PUBLIC -std=c++17) +endif() diff --git a/tests/catch/kernels/Set.cpp b/tests/catch/kernels/Set.cpp new file mode 100644 index 0000000000..72ad74dde2 --- /dev/null +++ b/tests/catch/kernels/Set.cpp @@ -0,0 +1,6 @@ +#include + +__global__ void Set(int* Ad, int val) { + int tx = threadIdx.x + blockIdx.x * blockDim.x; + Ad[tx] = val; +} \ No newline at end of file diff --git a/tests/catch/kernels/vectorADD.inl b/tests/catch/kernels/vectorADD.inl new file mode 100644 index 0000000000..4be89b6e89 --- /dev/null +++ b/tests/catch/kernels/vectorADD.inl @@ -0,0 +1,10 @@ +namespace HipTest { +template __global__ void vectorADD(const T* A_d, const T* B_d, T* C_d, size_t NELEM) { + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (size_t i = offset; i < NELEM; i += stride) { + C_d[i] = A_d[i] + B_d[i]; + } +} +} \ No newline at end of file diff --git a/tests/catch/unit/event/Unit_hipEventRecord.cc b/tests/catch/unit/event/Unit_hipEventRecord.cc index c3dc1a6086..d2f13d6e12 100644 --- a/tests/catch/unit/event/Unit_hipEventRecord.cc +++ b/tests/catch/unit/event/Unit_hipEventRecord.cc @@ -23,8 +23,8 @@ THE SOFTWARE. // Through manual inspection of the reported timestamps, can determine if recording a NULL event // forces synchronization : set #include -#include - +#include +#include #include TEST_CASE("Unit_hipEventRecord") { @@ -61,8 +61,8 @@ TEST_CASE("Unit_hipEventRecord") { // Record the start event HIP_CHECK(hipEventRecord(start, NULL)); - hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, - static_cast(A_d), static_cast(B_d), C_d, N); + HipTest::launchKernel(HipTest::vectorADD, blocks, threadsPerBlock, 0, 0, +static_cast(A_d), static_cast(B_d), C_d, N); HIP_CHECK(hipEventRecord(stop, NULL)); HIP_CHECK(hipEventSynchronize(stop)); @@ -87,5 +87,5 @@ TEST_CASE("Unit_hipEventRecord") { HIP_CHECK(hipEventDestroy(stop)); HipTest::checkVectorADD(A_h, B_h, C_h, N, true); - + TestContext::get().cleanContext(); } diff --git a/tests/catch/unit/memory/hipHostMalloc.cc b/tests/catch/unit/memory/hipHostMalloc.cc index 70459cf9d2..94f943c6bf 100644 --- a/tests/catch/unit/memory/hipHostMalloc.cc +++ b/tests/catch/unit/memory/hipHostMalloc.cc @@ -30,8 +30,9 @@ This testfile verifies the following scenarios of hipHostMalloc API */ #include -#include +#include #include +#include #define SYNC_EVENT 0 #define SYNC_STREAM 1 @@ -41,11 +42,6 @@ std::vector syncMsg = {"event", "stream", "device"}; static constexpr int numElements{1024 * 16}; static constexpr size_t sizeBytes{numElements * sizeof(int)}; -__global__ void Set(int* Ad, int val) { - int tx = threadIdx.x + blockIdx.x * blockDim.x; - Ad[tx] = val; -} - void CheckHostPointer(int numElements, int* ptr, unsigned eventFlags, int syncMethod, std::string msg) { std::cerr << "test: CheckHostPointer " @@ -70,10 +66,10 @@ void CheckHostPointer(int numElements, int* ptr, unsigned eventFlags, const int expected = 13; // Init array to know state: - hipLaunchKernelGGL(Set, dimGrid, dimBlock, 0, 0x0, ptr, -42); + HipTest::launchKernel(Set, dimGrid, dimBlock, 0, 0x0, ptr, -42); HIP_CHECK(hipDeviceSynchronize()); - hipLaunchKernelGGL(Set, dimGrid, dimBlock, 0, s, ptr, expected); + HipTest::launchKernel(Set, dimGrid, dimBlock, 0, s, ptr, expected); HIP_CHECK(hipEventRecord(e, s)); // Host waits for event : @@ -137,9 +133,9 @@ TEST_CASE("Unit_hipHostMalloc_Basic") { dim3 dimGrid(LEN / 512, 1, 1); dim3 dimBlock(512, 1, 1); - hipLaunchKernelGGL(HipTest::vectorADD, dimGrid, dimBlock, + HipTest::launchKernel(HipTest::vectorADD, dimGrid, dimBlock, 0, 0, static_cast(A_d), - static_cast(B_d), C_d, LEN); + static_cast(B_d), C_d, static_cast(LEN)); HIP_CHECK(hipMemcpy(C_h, C_d, LEN*sizeof(float), hipMemcpyDeviceToHost)); HIP_CHECK(hipDeviceSynchronize()); @@ -148,6 +144,7 @@ TEST_CASE("Unit_hipHostMalloc_Basic") { HIP_CHECK(hipHostFree(A_h)); HIP_CHECK(hipHostFree(B_h)); HIP_CHECK(hipHostFree(C_h)); + TestContext::get().cleanContext(); } } /* @@ -185,6 +182,7 @@ TEST_CASE("Unit_hipHostMalloc_NonCoherent") { SYNC_STREAM, ptrType); CheckHostPointer(numElements, A, hipEventReleaseToSystem, SYNC_EVENT, ptrType); + TestContext::get().cleanContext(); } /*