diff --git a/projects/hip-tests/catch/unit/rtc/CMakeLists.txt b/projects/hip-tests/catch/unit/rtc/CMakeLists.txt index c730ac82b4..39add7d88b 100644 --- a/projects/hip-tests/catch/unit/rtc/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/rtc/CMakeLists.txt @@ -15,6 +15,7 @@ set(AMD_TEST_SRC linker.cc shfl.cc stdheaders.cc + hiprtc_MathConstants_HeaderTst.cc ) add_custom_target(copyRtcHeaders ALL diff --git a/projects/hip-tests/catch/unit/rtc/hiprtc_MathConstants_HeaderTst.cc b/projects/hip-tests/catch/unit/rtc/hiprtc_MathConstants_HeaderTst.cc new file mode 100644 index 0000000000..2879c2e267 --- /dev/null +++ b/projects/hip-tests/catch/unit/rtc/hiprtc_MathConstants_HeaderTst.cc @@ -0,0 +1,231 @@ +/* +Copyright (c) 2023 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. +*/ + +/** +* @addtogroup hiprtc_MathConstants_HeaderTst hiprtc_MathConstants_HeaderTst +* @{ +* @ingroup hiprtcHeaders +* `hiprtcResult hiprtcCompileProgram(hiprtcProgram prog, +* int numOptions, +* const char** options);` - +* These test cases are target including various header file in kernel +* string and compile using the api mentioned above. +*/ + +#include +#include +#include + +static constexpr auto mathConstants_string { +R"( +extern "C" +__global__ void mathConstants(float *res) { + // single precision constants + res[0] = (HIP_INF_F == __int_as_float(0x7f800000U)); + res[1] = (HIP_REMQUO_MASK_F == (~((~0U)<= 6.0 +*/ +TEST_CASE("Unit_Rtc_MathConstants_header") { + std::string kernel_name = "mathConstants"; + const char* kername = kernel_name.c_str(); + float *result_h; + float *result_d; + int n = 93; + float Nbytes = n * sizeof(float); + result_h = new float[n]; + for (int i = 0; i < n; i++) { + result_h[i] = 0; + } + HIP_CHECK(hipMalloc(&result_d, Nbytes)); + HIP_CHECK(hipMemcpy(result_d, result_h, Nbytes, hipMemcpyHostToDevice)); + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + std::string architecture = prop.gcnArchName; + std::string complete_CO = "--gpu-architecture=" + architecture; + const char* compiler_option = complete_CO.c_str(); + for (int scenario = 0; scenario < 2; scenario++) { + hiprtcProgram prog; + const char** compiler_options = new const char*[2]; + for (int i = 0; i < 2; i++) { + compiler_options[0] = compiler_option; + compiler_options[1] = ""; + } + HIPRTC_CHECK(hiprtcCreateProgram(&prog, mathConstants_string, + kername, 0, NULL, NULL)); + if (scenario == 0) { + hiprtcResult compileResult{hiprtcCompileProgram(prog, + 1, &compiler_option)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("hiprtcCompileProgram() api failed!!"); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + REQUIRE(false); + } + } else { + hiprtcResult compileResult{hiprtcCompileProgram(prog, + 2, compiler_options)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("hiprtcCompileProgram() api failed!!"); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + REQUIRE(false); + } + } + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + std::vector codec(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, codec.data())); + void* kernelParam[] = {result_d}; + auto size = sizeof(kernelParam); + void* kernel_parameter[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &kernelParam, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + hipModule_t module; + hipFunction_t function; + HIP_CHECK(hipModuleLoadData(&module, codec.data())); + HIP_CHECK(hipModuleGetFunction(&function, module, kername)); + HIP_CHECK(hipModuleLaunchKernel(function, 1, 1, 1, 1, 1, 1, 0, 0, nullptr, + kernel_parameter)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(result_h, result_d, Nbytes, hipMemcpyDeviceToHost)); + for (int i = 0; i < n; i++) { + if (result_h[i] != 1.0f) { + WARN("FAIL for " << i << " iteration"); + WARN(result_h[i]); + REQUIRE(false); + } + } + HIP_CHECK(hipModuleUnload(module)); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + } + HIP_CHECK(hipFree(result_d)); + delete [] result_h; + REQUIRE(true); +}