From 90f5e7bf0c0a9e4908ed0e8ee8a741b1c7ddc21a Mon Sep 17 00:00:00 2001 From: Lakhan singh Thakur Date: Thu, 22 Feb 2024 11:37:44 +0530 Subject: [PATCH] SWDEV-379212 - Adding Tests for hiprtcGetLoweredName() api for templated kernels Change-Id: I45f092bb260dcd60709e79cbd7babfd2d9982858 [ROCm/hip-tests commit: 5f3c4b9ca1617fde63283c521230dd9661339788] --- .../hip-tests/catch/unit/rtc/CMakeLists.txt | 1 + .../catch/unit/rtc/hiprtcGetLoweredName.cc | 164 ++++++++++++++++++ 2 files changed, 165 insertions(+) create mode 100644 projects/hip-tests/catch/unit/rtc/hiprtcGetLoweredName.cc diff --git a/projects/hip-tests/catch/unit/rtc/CMakeLists.txt b/projects/hip-tests/catch/unit/rtc/CMakeLists.txt index e39d97b9b6..7a3bdd4ba9 100644 --- a/projects/hip-tests/catch/unit/rtc/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/rtc/CMakeLists.txt @@ -6,6 +6,7 @@ set(TEST_SRC hipStreamCaptureRtc.cc includepath.cc hipRTCDeviceMalloc.cc + hiprtcGetLoweredName.cc ) # AMD only tests diff --git a/projects/hip-tests/catch/unit/rtc/hiprtcGetLoweredName.cc b/projects/hip-tests/catch/unit/rtc/hiprtcGetLoweredName.cc new file mode 100644 index 0000000000..ce622a82d3 --- /dev/null +++ b/projects/hip-tests/catch/unit/rtc/hiprtcGetLoweredName.cc @@ -0,0 +1,164 @@ +/* +Copyright (c) 2024 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. +*/ + +#include +#include +#include +#include +#include + + +/** +* @addtogroup hiprtcGetLoweredName hiprtcGetLoweredName +* @{ +* @ingroup hiprtc +* `hiprtcResult hiprtcGetLoweredName(prog, +* kernel_name_vec[i].c_str(), // name expression +* &name // lowered name +* ));` - +* These test cases tests working hiprtcGetLoweredName() api +*/ + +static const char* const gpuProgram1 = R"( +template +__global__ void my_kernel(T* data) { + if (blockIdx.x != 0 || threadIdx.x != 0) return; + T data0 = data[0]; + for( int i=0; i +__global__ void my_kernel(T* data, K* Arr) { + if (blockIdx.x != 0 || threadIdx.x != 0) return; + T data0 = data[0]; + for( int i=0; i +__global__ void nontype_kernel() {} +)"; + +bool Test(int CaseNum, const char* GpuProgram) { + bool IfTestPassed = true; + // Create an instance of hiprtcProgram + hiprtcProgram prog; + + HIPRTC_CHECK(hiprtcCreateProgram(&prog, // prog + GpuProgram, // buffer + "prog.cu", // name + 0, // numHeaders + NULL, // headers + NULL)); // includeNames + + // add all name expressions for kernels + std::vector kernel_name_vec, kernelNameExpectdOutput; + if (CaseNum == 1) { + kernel_name_vec.push_back("my_kernel<(int)3, float >"); + kernel_name_vec.push_back("my_kernel<(int)5, double >"); + kernel_name_vec.push_back("my_kernel<(int)6, int >"); + kernel_name_vec.push_back("my_kernel<(int)10, long >"); + kernel_name_vec.push_back("my_kernel<(int)11, long long >"); + kernel_name_vec.push_back("my_kernel<(int)123, unsigned int >"); + kernel_name_vec.push_back("my_kernel<(int)1234, char >"); + kernel_name_vec.push_back("my_kernel<(int)12345, unsigned char >"); + + kernelNameExpectdOutput.push_back("_Z9my_kernelILi3EfEvPT0_"); + kernelNameExpectdOutput.push_back("_Z9my_kernelILi5EdEvPT0_"); + kernelNameExpectdOutput.push_back("_Z9my_kernelILi6EiEvPT0_"); + kernelNameExpectdOutput.push_back("_Z9my_kernelILi10ElEvPT0_"); + kernelNameExpectdOutput.push_back("_Z9my_kernelILi11ExEvPT0_"); + kernelNameExpectdOutput.push_back("_Z9my_kernelILi123EjEvPT0_"); + kernelNameExpectdOutput.push_back("_Z9my_kernelILi1234EcEvPT0_"); + kernelNameExpectdOutput.push_back("_Z9my_kernelILi12345EhEvPT0_"); + } else if (CaseNum == 2) { + kernel_name_vec.push_back("my_kernel<(int)3, float, float >"); + kernel_name_vec.push_back("my_kernel<(int)66, float, int >"); + kernel_name_vec.push_back("my_kernel<(int)(4 + 6), long, double >"); + kernel_name_vec.push_back("my_kernel<(int)5, double, double >"); + kernel_name_vec.push_back("my_kernel<(int)6, int, int >"); + kernel_name_vec.push_back("my_kernel<(int)10, long, long >"); + kernel_name_vec.push_back("my_kernel<(int)11, long long, long long >"); + kernel_name_vec.push_back("my_kernel<(int)123, unsigned int, unsigned int >"); //NOLINT + kernel_name_vec.push_back("my_kernel<(int)1234, char, char >"); + kernel_name_vec.push_back("my_kernel<(int)12345, unsigned char, unsigned char >"); //NOLINT + + kernelNameExpectdOutput.push_back("_Z9my_kernelILi3EffEvPT0_PT1_"); + kernelNameExpectdOutput.push_back("_Z9my_kernelILi66EfiEvPT0_PT1_"); + kernelNameExpectdOutput.push_back("_Z9my_kernelILi10EldEvPT0_PT1_"); + kernelNameExpectdOutput.push_back("_Z9my_kernelILi5EddEvPT0_PT1_"); + kernelNameExpectdOutput.push_back("_Z9my_kernelILi6EiiEvPT0_PT1_"); + kernelNameExpectdOutput.push_back("_Z9my_kernelILi10EllEvPT0_PT1_"); + kernelNameExpectdOutput.push_back("_Z9my_kernelILi11ExxEvPT0_PT1_"); + kernelNameExpectdOutput.push_back("_Z9my_kernelILi123EjjEvPT0_PT1_"); + kernelNameExpectdOutput.push_back("_Z9my_kernelILi1234EccEvPT0_PT1_"); + kernelNameExpectdOutput.push_back("_Z9my_kernelILi12345EhhEvPT0_PT1_"); + } else if (CaseNum == 3) { + kernel_name_vec.push_back("nontype_kernel<3>"); + kernel_name_vec.push_back("nontype_kernel<(2 + 3)>"); + kernel_name_vec.push_back("nontype_kernel<(-2 + 6)>"); + + kernelNameExpectdOutput.push_back("_Z14nontype_kernelILi3EEvv"); + kernelNameExpectdOutput.push_back("_Z14nontype_kernelILi5EEvv"); + kernelNameExpectdOutput.push_back("_Z14nontype_kernelILi4EEvv"); + } + + // add kernel name expressions to HIPRTC + for (size_t i = 0; i < kernel_name_vec.size(); ++i) + HIPRTC_CHECK(hiprtcAddNameExpression(prog, kernel_name_vec[i].c_str())); + + hiprtcResult compileResult = hiprtcCompileProgram(prog, // prog + 0, // numOptions + NULL); // options + // Obtain compilation log from the program. + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + + if (logSize > 0) { + char *log = new char[logSize]; + HIPRTC_CHECK(hiprtcGetProgramLog(prog, log)); + INFO(log); + delete[] log; + REQUIRE(compileResult == HIPRTC_SUCCESS); + } + // Extract lowered names + for (size_t i = 0; i < kernel_name_vec.size(); ++i) { + const char *name; + HIPRTC_CHECK(hiprtcGetLoweredName(prog, + kernel_name_vec[i].c_str(), // name expression + &name)); // lowered name + if (name != kernelNameExpectdOutput[i]) { + IfTestPassed = false; + } + } + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + return IfTestPassed; +} + + +TEST_CASE("Unit_hiprtcGetLoweredName_templateKrnls") { + REQUIRE(Test(1, gpuProgram1)); + REQUIRE(Test(2, gpuProgram2)); + REQUIRE(Test(3, gpuProgram3)); +}