From b8cc5e4a333363cda1205abdeba0927c06fa708e Mon Sep 17 00:00:00 2001 From: TomSang Date: Tue, 14 Sep 2021 04:07:56 -0400 Subject: [PATCH] SWDEV-299755 - Fix hiprtc failure on NV (#2341) 1. Add missing nvrtc lib. 2. Add context creation before call module(driver) api. 3. Remove obsolete saxpy_pch test. Change-Id: I5add27ebabcf021fa3ef99d09004c5d13423a297 [ROCm/hip commit: de8bb09c29d8d347048626b33f37499e467006d1] --- .../tests/src/hiprtc/hiprtcGetLoweredName.cpp | 20 ++- projects/hip/tests/src/hiprtc/saxpy.cpp | 29 +++- projects/hip/tests/src/hiprtc/saxpy_pch.cpp | 161 ------------------ 3 files changed, 37 insertions(+), 173 deletions(-) delete mode 100644 projects/hip/tests/src/hiprtc/saxpy_pch.cpp diff --git a/projects/hip/tests/src/hiprtc/hiprtcGetLoweredName.cpp b/projects/hip/tests/src/hiprtc/hiprtcGetLoweredName.cpp index 62ecf7beea..562f3aa2a8 100644 --- a/projects/hip/tests/src/hiprtc/hiprtcGetLoweredName.cpp +++ b/projects/hip/tests/src/hiprtc/hiprtcGetLoweredName.cpp @@ -20,7 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia + * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -lnvrtc * TEST: %t * HIT_END */ @@ -81,7 +81,12 @@ int main() hipDeviceProp_t props; int device = 0; hipGetDeviceProperties(&props, device); +#ifdef __HIP_PLATFORM_AMD__ std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName; +#else + std::string sarg = std::string("--gpu-architecture=compute_") + + std::to_string(props.major) + std::to_string(props.minor); +#endif const char* options[] = { sarg.c_str() }; @@ -108,11 +113,16 @@ int main() hiprtcGetCode(prog, code.data()); hipModule_t module; +#ifdef __HIP_PLATFORM_NVIDIA__ + HIPCHECK(hipInit(0)); + hipCtx_t ctx; + HIPCHECK(hipCtxCreate(&ctx, 0, device)); +#endif hipModuleLoadData(&module, code.data()); hipDeviceptr_t dResult; int hResult = 0; - hipMalloc(&dResult, sizeof(hResult)); + hipMalloc((void **)&dResult, sizeof(hResult)); hipMemcpyHtoD(dResult, &hResult, sizeof(hResult)); for (decltype(variable_name_vec.size()) i = 0; i != variable_name_vec.size(); ++i) { @@ -149,10 +159,12 @@ int main() if (expected_result[i] != hResult) { failed("Validation failed."); } } - hipFree(dResult); + hipFree((void *)dResult); hipModuleUnload(module); - hiprtcDestroyProgram(&prog); +#ifdef __HIP_PLATFORM_NVIDIA__ + HIPCHECK(hipCtxDestroy(ctx)); +#endif passed(); } diff --git a/projects/hip/tests/src/hiprtc/saxpy.cpp b/projects/hip/tests/src/hiprtc/saxpy.cpp index 16354eadba..f4acb961b1 100644 --- a/projects/hip/tests/src/hiprtc/saxpy.cpp +++ b/projects/hip/tests/src/hiprtc/saxpy.cpp @@ -20,11 +20,10 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia + * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -lnvrtc * TEST: %t * HIT_END */ - #include #include @@ -77,7 +76,12 @@ int main() hipDeviceProp_t props; int device = 0; hipGetDeviceProperties(&props, device); +#ifdef __HIP_PLATFORM_AMD__ std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName; +#else + std::string sarg = std::string("--gpu-architecture=compute_") + + std::to_string(props.major) + std::to_string(props.minor); +#endif const char* options[] = { sarg.c_str() }; @@ -106,6 +110,12 @@ int main() hipModule_t module; hipFunction_t kernel; + +#ifdef __HIP_PLATFORM_NVIDIA__ + HIPCHECK(hipInit(0)); + hipCtx_t ctx; + HIPCHECK(hipCtxCreate(&ctx, 0, device)); +#endif hipModuleLoadData(&module, code.data()); hipModuleGetFunction(&kernel, module, "saxpy"); @@ -123,9 +133,9 @@ int main() } hipDeviceptr_t dX, dY, dOut; - hipMalloc(&dX, bufferSize); - hipMalloc(&dY, bufferSize); - hipMalloc(&dOut, bufferSize); + hipMalloc((void **)&dX, bufferSize); + hipMalloc((void **)&dY, bufferSize); + hipMalloc((void **)&dOut, bufferSize); hipMemcpyHtoD(dX, hX.get(), bufferSize); hipMemcpyHtoD(dY, hY.get(), bufferSize); @@ -150,11 +160,14 @@ int main() if (fabs(a * hX[i] + hY[i] - hOut[i]) > fabs(hOut[i])* 1e-6) { failed("Validation failed."); } } - hipFree(dX); - hipFree(dY); - hipFree(dOut); + hipFree((void *)dX); + hipFree((void *)dY); + hipFree((void *)dOut); hipModuleUnload(module); +#ifdef __HIP_PLATFORM_NVIDIA__ + HIPCHECK(hipCtxDestroy(ctx)); +#endif passed(); } diff --git a/projects/hip/tests/src/hiprtc/saxpy_pch.cpp b/projects/hip/tests/src/hiprtc/saxpy_pch.cpp deleted file mode 100644 index 642b7c2971..0000000000 --- a/projects/hip/tests/src/hiprtc/saxpy_pch.cpp +++ /dev/null @@ -1,161 +0,0 @@ -/* -Copyright (c) 2015 - 2021 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. -*/ -/* HIT_START - * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia - * TEST: %t - * HIT_END - */ - -#include - -#include -#include - -#include -#include -#include -#include -#include -#include - -static constexpr auto NUM_THREADS{128}; -static constexpr auto NUM_BLOCKS{32}; - -static constexpr auto saxpy{ -R"( -#include "test_header.h" -#include "test_header1.h" -extern "C" -__global__ -void saxpy(real a, realptr x, realptr y, realptr out, size_t n) -{ - size_t tid = blockIdx.x * blockDim.x + threadIdx.x; - if (tid < n) { - out[tid] = a * x[tid] + y[tid] ; - } -} -)"}; - -int main() -{ - using namespace std; - - hiprtcProgram prog; - int num_headers = 2; - std::vector header_names; - std::vector header_sources; - header_names.push_back("test_header.h"); - header_names.push_back("test_header1.h"); - header_sources.push_back("#ifndef HIPRTC_TEST_HEADER_H\n#define HIPRTC_TEST_HEADER_H\ntypedef float real;\n#endif //HIPRTC_TEST_HEADER_H\n"); - header_sources.push_back("#ifndef HIPRTC_TEST_HEADER1_H\n#define HIPRTC_TEST_HEADER1_H\ntypedef float* realptr;\n#endif //HIPRTC_TEST_HEADER1_H\n"); - hiprtcCreateProgram(&prog, // prog - saxpy, // buffer - "saxpy.cu", // name - num_headers, // numHeaders - &header_sources[0], // headers - &header_names[0]); // includeNames - - hipDeviceProp_t props; - int device = 0; - hipGetDeviceProperties(&props, device); - std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName; - const char* options[] = { - "-hip-pch", - sarg.c_str() - }; - - hiprtcResult compileResult{hiprtcCompileProgram(prog, 2, options)}; - - size_t logSize; - hiprtcGetProgramLogSize(prog, &logSize); - - if (logSize) { - string log(logSize, '\0'); - hiprtcGetProgramLog(prog, &log[0]); - - cout << log << '\n'; - } - - if (compileResult != HIPRTC_SUCCESS) { failed("Compilation failed."); } - - size_t codeSize; - hiprtcGetCodeSize(prog, &codeSize); - - vector code(codeSize); - hiprtcGetCode(prog, code.data()); - - hiprtcDestroyProgram(&prog); - - hipModule_t module; - hipFunction_t kernel; - hipModuleLoadData(&module, code.data()); - hipModuleGetFunction(&kernel, module, "saxpy"); - - size_t n = NUM_THREADS * NUM_BLOCKS; - size_t bufferSize = n * sizeof(float); - - float a = 5.1f; - unique_ptr hX{new float[n]}; - unique_ptr hY{new float[n]}; - unique_ptr hOut{new float[n]}; - - for (size_t i = 0; i < n; ++i) { - hX[i] = static_cast(i); - hY[i] = static_cast(i * 2); - } - - hipDeviceptr_t dX, dY, dOut; - hipMalloc(&dX, bufferSize); - hipMalloc(&dY, bufferSize); - hipMalloc(&dOut, bufferSize); - hipMemcpyHtoD(dX, hX.get(), bufferSize); - hipMemcpyHtoD(dY, hY.get(), bufferSize); - - struct { - float a_; - hipDeviceptr_t b_; - hipDeviceptr_t c_; - hipDeviceptr_t d_; - size_t e_; - } args{a, dX, dY, dOut, n}; - - auto size = sizeof(args); - void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, - HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, - HIP_LAUNCH_PARAM_END}; - - hipModuleLaunchKernel(kernel, NUM_BLOCKS, 1, 1, NUM_THREADS, 1, 1, - 0, nullptr, nullptr, config); - hipMemcpyDtoH(hOut.get(), dOut, bufferSize); - - for (size_t i = 0; i < n; ++i) { - if (fabs(a * hX[i] + hY[i] - hOut[i]) > fabs(hOut[i])* 1e-6) { failed("Validation failed."); } - } - - hipFree(dX); - hipFree(dY); - hipFree(dOut); - - hipModuleUnload(module); - - passed(); -}