From 4c31bfcc612def8ad3e83df2b376936767a50e1d Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Tue, 20 Jul 2021 04:13:19 -0700 Subject: [PATCH] SWDEV-289409 - Add hiprtc tests Change-Id: Ib26527b704aed32ae3f3ed38bf6e2fd412462c8e [ROCm/hip-tests commit: 46fb008ba64adff860e0f0a318b028b6ab9fb87c] --- .../catch/hipTestMain/CMakeLists.txt | 6 +- .../catch/include/hip_test_checkers.hh | 14 ++-- .../catch/include/hip_test_common.hh | 9 +++ projects/hip-tests/catch/unit/CMakeLists.txt | 4 +- .../hip-tests/catch/unit/rtc/CMakeLists.txt | 12 ++- projects/hip-tests/catch/unit/rtc/saxpy.cc | 73 ++++++++++--------- 6 files changed, 64 insertions(+), 54 deletions(-) diff --git a/projects/hip-tests/catch/hipTestMain/CMakeLists.txt b/projects/hip-tests/catch/hipTestMain/CMakeLists.txt index 449793b630..c19f48aa85 100644 --- a/projects/hip-tests/catch/hipTestMain/CMakeLists.txt +++ b/projects/hip-tests/catch/hipTestMain/CMakeLists.txt @@ -14,11 +14,11 @@ target_link_libraries(UnitTests PRIVATE UnitDeviceTests StreamTest OccupancyTest DeviceTest + RTC stdc++fs) -# Add AMD Only Tests -if(HIP_PLATFORM MATCHES "amd") - # target_link_libraries(UnitTests PRIVATE RTC) +if(HIP_PLATFORM MATCHES "nvidia") + target_link_libraries(UnitTests PRIVATE nvrtc) endif() catch_discover_tests(UnitTests PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST") diff --git a/projects/hip-tests/catch/include/hip_test_checkers.hh b/projects/hip-tests/catch/include/hip_test_checkers.hh index a96e597b37..e3941d17c6 100644 --- a/projects/hip-tests/catch/include/hip_test_checkers.hh +++ b/projects/hip-tests/catch/include/hip_test_checkers.hh @@ -2,13 +2,13 @@ #include "hip_test_common.hh" #include using namespace std; -#define guarantee(cond, str) \ - { \ - if (!(cond)) { \ - std::cout << str << std::endl; \ - abort(); \ - } \ - } +#define guarantee(cond, str) \ + { \ + if (!(cond)) { \ + std::cout << str << std::endl; \ + abort(); \ + } \ + } namespace HipTest { diff --git a/projects/hip-tests/catch/include/hip_test_common.hh b/projects/hip-tests/catch/include/hip_test_common.hh index e8b2f40f88..d5041be9b9 100644 --- a/projects/hip-tests/catch/include/hip_test_common.hh +++ b/projects/hip-tests/catch/include/hip_test_common.hh @@ -36,6 +36,15 @@ THE SOFTWARE. } \ } +#define HIPRTC_CHECK(error) \ + { \ + auto localError = error; \ + if (localError != HIPRTC_SUCCESS) { \ + INFO("Error: " << hiprtcGetErrorString(localError) << " Code: " << localError << " Str: " \ + << #error << " In File: " << __FILE__ << " At line: " << __LINE__); \ + REQUIRE(false); \ + } \ + } // Although its assert, it will be evaluated at runtime #define HIP_ASSERT(x) \ { REQUIRE((x)); } diff --git a/projects/hip-tests/catch/unit/CMakeLists.txt b/projects/hip-tests/catch/unit/CMakeLists.txt index d5676d59e9..be8a4f2d20 100644 --- a/projects/hip-tests/catch/unit/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/CMakeLists.txt @@ -3,6 +3,4 @@ add_subdirectory(deviceLib) add_subdirectory(stream) add_subdirectory(occupancy) add_subdirectory(device) - -# Disable Saxpy test temporarily to see if CI Passes -# add_subdirectory(rtc) +add_subdirectory(rtc) diff --git a/projects/hip-tests/catch/unit/rtc/CMakeLists.txt b/projects/hip-tests/catch/unit/rtc/CMakeLists.txt index 062e4153c6..d52d3f9288 100644 --- a/projects/hip-tests/catch/unit/rtc/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/rtc/CMakeLists.txt @@ -1,12 +1,10 @@ # AMD Tests -set(AMD_TEST_SRC +set(TEST_SRC saxpy.cc ) -if(HIP_PLATFORM MATCHES "amd") - # Create shared lib of all tests - add_library(RTC SHARED EXCLUDE_FROM_ALL ${AMD_TEST_SRC}) +# Create shared lib of all tests +add_library(RTC SHARED EXCLUDE_FROM_ALL ${TEST_SRC}) - # Add dependency on build_tests to build it on this custom target - add_dependencies(build_tests RTC) -endif() +# Add dependency on build_tests to build it on this custom target +add_dependencies(build_tests RTC) diff --git a/projects/hip-tests/catch/unit/rtc/saxpy.cc b/projects/hip-tests/catch/unit/rtc/saxpy.cc index 09114a4453..a200b30d19 100644 --- a/projects/hip-tests/catch/unit/rtc/saxpy.cc +++ b/projects/hip-tests/catch/unit/rtc/saxpy.cc @@ -15,20 +15,19 @@ static constexpr auto NUM_THREADS{128}; static constexpr auto NUM_BLOCKS{32}; static constexpr auto saxpy{ -R"( + R"( extern "C" __global__ void saxpy(float a, float* x, float* y, float* out, size_t n) { size_t tid = blockIdx.x * blockDim.x + threadIdx.x; - if (tid < n) { - out[tid] = a * x[tid] + y[tid]; - } - + if (tid < n) { + out[tid] = a * x[tid] + y[tid]; + } } )"}; -TEST_CASE("saxpy", "[hiprtc][saxpy]") { +TEST_CASE("Unit_hiprtc_saxpy") { using namespace std; hiprtcProgram prog; hiprtcCreateProgram(&prog, // prog @@ -37,34 +36,43 @@ TEST_CASE("saxpy", "[hiprtc][saxpy]") { 0, nullptr, nullptr); hipDeviceProp_t props; int device = 0; - hipGetDeviceProperties(&props, device); + HIP_CHECK(hipGetDeviceProperties(&props, device)); +#ifdef __HIP_PLATFORM_AMD__ std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName; +#else + std::string sarg = std::string("--fmad=false"); +#endif const char* options[] = {sarg.c_str()}; hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)}; size_t logSize; - hiprtcGetProgramLogSize(prog, &logSize); + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); if (logSize) { string log(logSize, '\0'); - hiprtcGetProgramLog(prog, &log[0]); - + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); std::cout << log << '\n'; } REQUIRE(compileResult == HIPRTC_SUCCESS); size_t codeSize; - hiprtcGetCodeSize(prog, &codeSize); + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); vector code(codeSize); - hiprtcGetCode(prog, code.data()); + HIPRTC_CHECK(hiprtcGetCode(prog, code.data())); - hiprtcDestroyProgram(&prog); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + + // Do hip malloc first so that we donot need to do a cuInit manually before calling hipModule APIs + size_t n = NUM_THREADS * NUM_BLOCKS; + size_t bufferSize = n * sizeof(float); + + float *dX, *dY, *dOut; + HIP_CHECK(hipMalloc(&dX, bufferSize)); + HIP_CHECK(hipMalloc(&dY, bufferSize)); + HIP_CHECK(hipMalloc(&dOut, bufferSize)); 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); + HIP_CHECK(hipModuleLoadData(&module, code.data())); + HIP_CHECK(hipModuleGetFunction(&kernel, module, "saxpy")); float a = 5.1f; unique_ptr hX{new float[n]}; @@ -75,18 +83,14 @@ TEST_CASE("saxpy", "[hiprtc][saxpy]") { 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); + HIP_CHECK(hipMemcpy(dX, hX.get(), bufferSize, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(dY, hY.get(), bufferSize, hipMemcpyHostToDevice)); struct { float a_; - hipDeviceptr_t b_; - hipDeviceptr_t c_; - hipDeviceptr_t d_; + float* b_; + float* c_; + float* d_; size_t e_; } args{a, dX, dY, dOut, n}; @@ -95,17 +99,18 @@ TEST_CASE("saxpy", "[hiprtc][saxpy]") { HIP_LAUNCH_PARAM_END}; hipModuleLaunchKernel(kernel, NUM_BLOCKS, 1, 1, NUM_THREADS, 1, 1, 0, nullptr, nullptr, config); - hipMemcpyDtoH(hOut.get(), dOut, bufferSize); + + HIP_CHECK(hipMemcpy(hOut.get(), dOut, bufferSize, hipMemcpyDeviceToHost)); + + hipFree(dX); + hipFree(dY); + hipFree(dOut); + + HIP_CHECK(hipModuleUnload(module)); for (size_t i = 0; i < n; ++i) { INFO("For " << i << " Value: " << fabs(a * hX[i] + hY[i] - hOut[i]) << " with: " << (fabs(hOut[i] * 1.0f) * 1e-6)); REQUIRE(fabs(a * hX[i] + hY[i] - hOut[i]) <= fabs(hOut[i]) * 1e-6); } - - hipFree(dX); - hipFree(dY); - hipFree(dOut); - - hipModuleUnload(module); }