SWDEV-254185 - Added support to pass include headers to hipRTC

Change-Id: Ia19e6806a476d3c6ecd3e9a89c3ed72953a4900c


[ROCm/clr commit: 7d67573a4f]

projects/clr/hipamd/rocclr/hip_rtc.cpp

@@ -41,9 +41,6 @@ public:
   std::map<std::string, std::pair<std::string, std::string>> 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<std::string> vHeaderNames;
-  std::vector<std::string> 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<char*>(it->second.first.data());
-    *headerNames = reinterpret_cast<char*>(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<hiprtcProgram>(as_cl(program));
 
   HIPRTC_RETURN(HIPRTC_SUCCESS);

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()
     };

projects/clr/hipamd/tests/src/hiprtc/saxpy.cpp

@@ -43,14 +43,15 @@ static constexpr auto NUM_BLOCKS{32};
 static constexpr auto saxpy{
 R"(
 #include <hip/hip_runtime.h>
-
+#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<const char*> header_names;
+    std::vector<const char*> 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) {