SWDEV-409299 - Vega clinfo is not working

Change-Id: Ia48bc6f130bd102dff210b105de6f9c02ebbe012


[ROCm/clr commit: 7912f3af89]

projects/clr/rocclr/device/blitcl.cpp

@@ -95,7 +95,9 @@ const char* BlitLinearSourceCode = BLIT_KERNELS(
                                                    ulong4 srcRect, ulong4 dstRect, ulong4 size) {
     __amd_copyBufferRectAligned(src, dst, srcRect, dstRect, size);
   }
+);
 
+const char* HipExtraSourceCode = BLIT_KERNELS(
   __kernel void __amd_rocclr_streamOpsWrite(__global uint* ptrInt, __global ulong* ptrUlong,
                                             ulong value, ulong sizeBytes) {
     __amd_streamOpsWrite(ptrInt, ptrUlong, value, sizeBytes);

projects/clr/rocclr/device/pal/paldevice.cpp

@@ -145,6 +145,10 @@ static std::tuple<Pal::GfxIpLevel, Pal::AsicRevision, const char*> findPal(uint3
 
 }  // namespace
 
+namespace device {
+extern const char* HipExtraSourceCode;
+}
+
 bool PalDeviceLoad() {
   bool ret = false;
 
@@ -2508,7 +2512,9 @@ bool Device::createBlitProgram() {
   // Delayed compilation due to brig_loader memory allocation
   std::string extraBlits;
   std::string ocl20;
-  if (!amd::IS_HIP) {
+  if (amd::IS_HIP) {
+    extraBlits = device::HipExtraSourceCode;
+  } else {
     if (settings().oclVersion_ >= OpenCL20) {
       extraBlits = iDev()->GetDispatchKernelSource();
       if (settings().useLightning_) {

projects/clr/rocclr/device/rocm/rocdevice.cpp

@@ -65,6 +65,7 @@
 #define OPENCL_C_VERSION_STR XSTR(OPENCL_C_MAJOR) "." XSTR(OPENCL_C_MINOR)
 
 #ifndef WITHOUT_HSA_BACKEND
+
 namespace {
 
 inline bool getIsaMeta(std::string isaName, amd_comgr_metadata_node_t& isaMeta) {
@@ -94,7 +95,7 @@ bool getValueFromIsaMeta(amd_comgr_metadata_node_t& isaMeta, const char* key,
 } // namespace
 
 namespace device {
-extern const char* BlitSourceCode;
+extern const char* HipExtraSourceCode;
 } // namespace device
 
 namespace roc {
@@ -843,7 +844,9 @@ bool Device::createBlitProgram() {
 
 #if defined(USE_COMGR_LIBRARY)
   if (settings().useLightning_) {
-    if (!amd::IS_HIP) {
+    if (amd::IS_HIP) {
+      extraKernel = device::HipExtraSourceCode;
+    } else {
       extraKernel = SchedulerSourceCode;
     }