From a7418845f8c65ead6592298e576b1684b650f61d Mon Sep 17 00:00:00 2001
From: Alex Xie <AlexBin.Xie@amd.com>
Date: Thu, 6 Jul 2023 16:13:14 -0400
Subject: [PATCH] SWDEV-409299 - Vega clinfo is not working

Change-Id: Ia48bc6f130bd102dff210b105de6f9c02ebbe012


[ROCm/clr commit: 7912f3af89e1afdb856da24f5ebc9f60283951e0]
---
 projects/clr/rocclr/device/blitcl.cpp         | 2 ++
 projects/clr/rocclr/device/pal/paldevice.cpp  | 8 +++++++-
 projects/clr/rocclr/device/rocm/rocdevice.cpp | 7 +++++--
 3 files changed, 14 insertions(+), 3 deletions(-)

diff --git a/projects/clr/rocclr/device/blitcl.cpp b/projects/clr/rocclr/device/blitcl.cpp
index cfae2b9430..6d6a507092 100644
--- a/projects/clr/rocclr/device/blitcl.cpp
+++ b/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);
diff --git a/projects/clr/rocclr/device/pal/paldevice.cpp b/projects/clr/rocclr/device/pal/paldevice.cpp
index 4192752f34..bc7f5ef3f3 100644
--- a/projects/clr/rocclr/device/pal/paldevice.cpp
+++ b/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_) {
diff --git a/projects/clr/rocclr/device/rocm/rocdevice.cpp b/projects/clr/rocclr/device/rocm/rocdevice.cpp
index 5acaab4a0e..8897edafff 100644
--- a/projects/clr/rocclr/device/rocm/rocdevice.cpp
+++ b/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;
     }