diff --git a/projects/clr/hipamd/rocclr/hip_rtc.cpp b/projects/clr/hipamd/rocclr/hip_rtc.cpp index 08ee8b48ce..ee66cb8f24 100755 --- a/projects/clr/hipamd/rocclr/hip_rtc.cpp +++ b/projects/clr/hipamd/rocclr/hip_rtc.cpp @@ -141,6 +141,34 @@ static std::string handleMangledName(std::string name) { return loweredName; } +static std::string getValueOf(const std::string& option) { + std::string res; + auto f = std::find(option.begin(), option.end(), '='); + if (f != option.end()) res = std::string(f + 1, option.end()); + return res; +} + +static void transformOptions(std::vector& options, amd::Program* program) { + std::vector t_option; + for (auto& i : options) { + std::string t_str(i); + // Use precompiled header for hip + if (t_str == "-hip-pch") { + const char* pch = nullptr; + unsigned int pch_size = 0; + __hipGetPCH(&pch, &pch_size); + program->addPreCompiledHeader(std::string(pch, pch_size)); + i = "-nogpuinc"; + } + // To maintain compatibility with nvrtc + if (t_str.find("--gpu-architecture=") != std::string::npos) { + auto value = getValueOf(t_str); + std::string newValue = "--offload-arch=" + value; + i = newValue.c_str(); + } + } +} + const char* hiprtcGetErrorString(hiprtcResult x) { switch (x) { case HIPRTC_SUCCESS: @@ -217,6 +245,7 @@ hiprtcResult hiprtcCompileProgram(hiprtcProgram prog, int numOptions, const char std::ostringstream ostrstr; std::vector oarr(&options[0], &options[numOptions]); + transformOptions(oarr, program); std::copy(oarr.begin(), oarr.end(), std::ostream_iterator(ostrstr, " ")); ostrstr.str().append(" -DHIP_VERSION_MAJOR=9"); diff --git a/projects/clr/hipamd/tests/src/hiprtc/saxpy_pch.cpp b/projects/clr/hipamd/tests/src/hiprtc/saxpy_pch.cpp new file mode 100755 index 0000000000..1fc403beeb --- /dev/null +++ b/projects/clr/hipamd/tests/src/hiprtc/saxpy_pch.cpp @@ -0,0 +1,161 @@ +/* +Copyright (c) 2015 - present 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 LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_RUNTIME rocclr + * 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[] = { + "-use-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(); +}