diff --git a/projects/clr/hipamd/rocclr/hip_rtc.cpp b/projects/clr/hipamd/rocclr/hip_rtc.cpp index 804e903893..08ee8b48ce 100755 --- a/projects/clr/hipamd/rocclr/hip_rtc.cpp +++ b/projects/clr/hipamd/rocclr/hip_rtc.cpp @@ -41,9 +41,6 @@ public: std::map> nameExpresssion_; static ProgramState& instance(); - void createProgramHeaders(amd::Program* program, int numHeaders, - const char** headers, const char** headerNames); - void getProgramHeaders(amd::Program* program, int* numHeaders, char** headers, char ** headerNames); uint32_t addNameExpression(const char* name_expression); char* getLoweredName(const char* name_expression); }; @@ -57,30 +54,6 @@ ProgramState& ProgramState::instance() { return *programState_; } -void ProgramState::createProgramHeaders(amd::Program* program, int numHeaders, - const char** headers, const char** headerNames) { - amd::ScopedLock lock(lock_); - std::vector vHeaderNames; - std::vector vHeaders; - for (auto i = 0; i != numHeaders; ++i) { - vHeaders.emplace_back(headers[i]); - vHeaderNames.emplace_back(headerNames[i]); - progHeaders_[program] = std::make_pair(std::move(vHeaders), std::move(vHeaderNames)); - } -} - -void ProgramState::getProgramHeaders(amd::Program* program, int* numHeaders, - char** headers, char ** headerNames) { - amd::ScopedLock lock(lock_); - - const auto it = progHeaders_.find(program); - if (it != progHeaders_.cend()) { - *numHeaders = it->second.first.size(); - *headers = reinterpret_cast(it->second.first.data()); - *headerNames = reinterpret_cast(it->second.second.data()); - } -} - uint32_t ProgramState::addNameExpression(const char* name_expression) { amd::ScopedLock lock(lock_); @@ -218,7 +191,8 @@ hiprtcResult hiprtcCreateProgram(hiprtcProgram* prog, const char* src, const cha HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); } - amd::Program* program = new amd::Program(*hip::getCurrentDevice()->asContext(), src, amd::Program::HIP); + amd::Program* program = new amd::Program(*hip::getCurrentDevice()->asContext(), src, amd::Program::HIP, + numHeaders, headers, headerNames); if (program == NULL) { HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); } @@ -228,8 +202,6 @@ hiprtcResult hiprtcCreateProgram(hiprtcProgram* prog, const char* src, const cha HIPRTC_RETURN(HIPRTC_ERROR_PROGRAM_CREATION_FAILURE); } - ProgramState::instance().createProgramHeaders(program, numHeaders, headers, headerNames); - *prog = reinterpret_cast(as_cl(program)); HIPRTC_RETURN(HIPRTC_SUCCESS); diff --git a/projects/clr/hipamd/tests/src/hiprtc/hiprtcGetLoweredName.cpp b/projects/clr/hipamd/tests/src/hiprtc/hiprtcGetLoweredName.cpp index 11c4ab1dba..639b790fd7 100644 --- a/projects/clr/hipamd/tests/src/hiprtc/hiprtcGetLoweredName.cpp +++ b/projects/clr/hipamd/tests/src/hiprtc/hiprtcGetLoweredName.cpp @@ -83,7 +83,7 @@ int main() hipDeviceProp_t props; int device = 0; hipGetDeviceProperties(&props, device); - std::string sarg = "--gpu-architecture=" + props.gcnArchName; + std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName; const char* options[] = { sarg.c_str() }; diff --git a/projects/clr/hipamd/tests/src/hiprtc/saxpy.cpp b/projects/clr/hipamd/tests/src/hiprtc/saxpy.cpp index f6489c761e..5090d645ca 100755 --- a/projects/clr/hipamd/tests/src/hiprtc/saxpy.cpp +++ b/projects/clr/hipamd/tests/src/hiprtc/saxpy.cpp @@ -43,14 +43,15 @@ static constexpr auto NUM_BLOCKS{32}; static constexpr auto saxpy{ R"( #include - +#include "test_header.h" +#include "test_header1.h" extern "C" __global__ -void saxpy(float a, float* x, float* y, float* out, size_t n) +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]; + out[tid] = a * x[tid] + y[tid] ; } } )"}; @@ -60,17 +61,24 @@ 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 - 0, // numHeaders - nullptr, // headers - nullptr); // includeNames + num_headers, // numHeaders + &header_sources[0], // headers + &header_names[0]); // includeNames hipDeviceProp_t props; int device = 0; hipGetDeviceProperties(&props, device); - std::string sarg = "--gpu-architecture=" + props.gcnArchName; + std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName; const char* options[] = { sarg.c_str() }; @@ -99,7 +107,6 @@ int main() hipModule_t module; hipFunction_t kernel; - hipModuleLoadData(&module, code.data()); hipModuleGetFunction(&kernel, module, "saxpy"); @@ -138,7 +145,6 @@ int main() 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) {