From de14b1ce6dfc3cf8feea1a55e91c5efd49ba103f Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 23 Sep 2022 12:36:57 +0530 Subject: [PATCH] SWDEV-344342 - Add sample using hiprtc target (#2946) Change-Id: Ie2a0b308862ac8db2e72a0170499139c7b46605b --- .../2_Cookbook/23_cmake_hiprtc/CMakeLists.txt | 36 ++++ samples/2_Cookbook/23_cmake_hiprtc/README.md | 9 + samples/2_Cookbook/23_cmake_hiprtc/saxpy.cpp | 156 ++++++++++++++++++ 3 files changed, 201 insertions(+) create mode 100644 samples/2_Cookbook/23_cmake_hiprtc/CMakeLists.txt create mode 100644 samples/2_Cookbook/23_cmake_hiprtc/README.md create mode 100644 samples/2_Cookbook/23_cmake_hiprtc/saxpy.cpp diff --git a/samples/2_Cookbook/23_cmake_hiprtc/CMakeLists.txt b/samples/2_Cookbook/23_cmake_hiprtc/CMakeLists.txt new file mode 100644 index 0000000000..f22b063bb3 --- /dev/null +++ b/samples/2_Cookbook/23_cmake_hiprtc/CMakeLists.txt @@ -0,0 +1,36 @@ +# 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. + +project(cmake_hiprtc_test) + +cmake_minimum_required(VERSION 3.10.2) + +# Find hiprtc +find_package(hiprtc REQUIRED) +# Find hip +find_package(hip REQUIRED) + +# Create the excutable +add_executable(test saxpy.cpp) + +# Link with HIPRTC +target_link_libraries(test hiprtc::hiprtc) +# Link with HIP +target_link_libraries(test hip::device) diff --git a/samples/2_Cookbook/23_cmake_hiprtc/README.md b/samples/2_Cookbook/23_cmake_hiprtc/README.md new file mode 100644 index 0000000000..2831619d19 --- /dev/null +++ b/samples/2_Cookbook/23_cmake_hiprtc/README.md @@ -0,0 +1,9 @@ +### This will test linking hiprtc::hiprtc interface in cmake +I. Build +mkdir -p build; cd build +rm -rf *; CXX=amdclang++ cmake -DCMAKE_PREFIX_PATH=/opt/rocm/hip .. +make + +II. Test +$ ./test +SAXPY test completed diff --git a/samples/2_Cookbook/23_cmake_hiprtc/saxpy.cpp b/samples/2_Cookbook/23_cmake_hiprtc/saxpy.cpp new file mode 100644 index 0000000000..9c24091f0f --- /dev/null +++ b/samples/2_Cookbook/23_cmake_hiprtc/saxpy.cpp @@ -0,0 +1,156 @@ +/* +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. +*/ + +#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; + vector header_names; + 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); + + const char* options[] = {}; + + hiprtcResult compileResult{hiprtcCompileProgram(prog, 0, options)}; + + size_t logSize; + hiprtcGetProgramLogSize(prog, &logSize); + + if (logSize) { + string log(logSize, '\0'); + hiprtcGetProgramLog(prog, &log[0]); + + cout << log << '\n'; + } + + if (compileResult != HIPRTC_SUCCESS) { + cout << "Compilation failed." << endl; + } + + 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((void **)&dX, bufferSize); + hipMalloc((void **)&dY, bufferSize); + hipMalloc((void **)&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) { + cout << "Validation failed." << endl; + } + } + + hipFree((void *)dX); + hipFree((void *)dY); + hipFree((void *)dOut); + + hipModuleUnload(module); + + cout << "SAXPY test completed" << endl; +}