diff --git a/catch/CMakeLists.txt b/catch/CMakeLists.txt index 40a61fd6f7..1a272d03aa 100644 --- a/catch/CMakeLists.txt +++ b/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/catch/external/Catch2/cmake/Catch2/Catch.cmake b/catch/external/Catch2/cmake/Catch2/Catch.cmake index 5287241143..62e4984afc 100644 --- a/catch/external/Catch2/cmake/Catch2/Catch.cmake +++ b/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/catch/hipTestMain/hip_test_context.cc b/catch/hipTestMain/hip_test_context.cc index ec55e2395e..a6e3a08609 100644 --- a/catch/hipTestMain/hip_test_context.cc +++ b/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/catch/include/hip_test_common.hh b/catch/include/hip_test_common.hh index 496f202b8c..63b8354421 100644 --- a/catch/include/hip_test_common.hh +++ b/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/catch/include/hip_test_context.hh b/catch/include/hip_test_context.hh index e6429454ae..2d630b70f7 100644 --- a/catch/include/hip_test_context.hh +++ b/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/catch/include/hip_test_rtc.hh b/catch/include/hip_test_rtc.hh new file mode 100644 index 0000000000..0860ae51db --- /dev/null +++ b/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/catch/include/kernel_mapping.hh b/catch/include/kernel_mapping.hh new file mode 100644 index 0000000000..32ca93b5e6 --- /dev/null +++ b/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/catch/include/kernels.hh b/catch/include/kernels.hh new file mode 100644 index 0000000000..ab2a92bfdb --- /dev/null +++ b/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/catch/kernels/CMakeLists.txt b/catch/kernels/CMakeLists.txt new file mode 100644 index 0000000000..91e1ab69ff --- /dev/null +++ b/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/catch/kernels/Set.cpp b/catch/kernels/Set.cpp new file mode 100644 index 0000000000..72ad74dde2 --- /dev/null +++ b/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/catch/kernels/vectorADD.inl b/catch/kernels/vectorADD.inl new file mode 100644 index 0000000000..4be89b6e89 --- /dev/null +++ b/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/catch/unit/event/Unit_hipEventRecord.cc b/catch/unit/event/Unit_hipEventRecord.cc index c3dc1a6086..d2f13d6e12 100644 --- a/catch/unit/event/Unit_hipEventRecord.cc +++ b/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/catch/unit/memory/hipHostMalloc.cc b/catch/unit/memory/hipHostMalloc.cc index 70459cf9d2..94f943c6bf 100644 --- a/catch/unit/memory/hipHostMalloc.cc +++ b/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(); } /*