From a932af9dbc172d5f4136524a1ddf7dc5bc021440 Mon Sep 17 00:00:00 2001
From: foreman <dl.swbuild@amd.com>
Date: Fri, 27 Sep 2019 14:07:19 -0400
Subject: [PATCH] P4 to Git Change 2005601 by gandryey@gera-hip-lnx on
 2019/09/27 13:58:55

	SWDEV-184709 - support hipLaunchCooperativeKernel()
	- Switch to the device library functions for GWS sync and init

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#42 edit
... //depot/stg/opencl/drivers/opencl/make/hip.git/tests/src/runtimeApi/module/hipCooperativeGroup.cpp#2 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palschedcl.cpp#7 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.cpp#152 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocschedcl.cpp#4 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#87 edit
---
 rocclr/runtime/device/pal/palschedcl.cpp  | 17 +++--------------
 rocclr/runtime/device/pal/palvirtual.cpp  |  5 +----
 rocclr/runtime/device/rocm/rocschedcl.cpp | 18 ++++--------------
 rocclr/runtime/device/rocm/rocvirtual.cpp |  5 +----
 4 files changed, 9 insertions(+), 36 deletions(-)

diff --git a/rocclr/runtime/device/pal/palschedcl.cpp b/rocclr/runtime/device/pal/palschedcl.cpp
index 159c564b95..bdbd26cb1e 100644
--- a/rocclr/runtime/device/pal/palschedcl.cpp
+++ b/rocclr/runtime/device/pal/palschedcl.cpp
@@ -17,21 +17,10 @@ __kernel void scheduler(__global void* queue, __global void* params, uint paramI
 
 const char* GwsInitSourceCode = BLIT_KERNEL(
 \n
+extern void __ockl_gws_init(uint nwm1, uint rid);
+\n
 __kernel void gwsInit(uint value) {
-    unsigned int m0_backup, new_m0;
-    __asm__ __volatile__(
-        "s_mov_b32 %0 m0\n"
-        "v_readfirstlane_b32 %1 %2\n"
-        "s_nop 0\n"
-        "s_mov_b32 m0 %1\n"
-        "s_nop 0\n"
-        "ds_gws_init %3 offset:0 gds\n"
-        "s_waitcnt lgkmcnt(0) expcnt(0)\n"
-        "s_mov_b32 m0 %0\n"
-        "s_nop 0"
-        : "=s"(m0_backup), "=s"(new_m0)
-        : "v"(0 << 0x10), "{v0}"(value - 1)
-        : "memory");
+  __ockl_gws_init(value, 0);
 }
 \n);
 
diff --git a/rocclr/runtime/device/pal/palvirtual.cpp b/rocclr/runtime/device/pal/palvirtual.cpp
index 78031b8ea1..1d4a9eb85a 100644
--- a/rocclr/runtime/device/pal/palvirtual.cpp
+++ b/rocclr/runtime/device/pal/palvirtual.cpp
@@ -2295,9 +2295,6 @@ void VirtualGPU::submitKernel(amd::NDRangeKernelCommand& vcmd) {
         workgroups += (vcmd.sizes().global()[i] / vcmd.sizes().local()[i]);
       }
     }
-    uint32_t counter = workgroups *
-        amd::alignUp(vcmd.sizes().local().product(), dev().info().wavefrontWidth_) /
-        dev().info().wavefrontWidth_;
 
     bool test = true;
     VirtualGPU* queue = (test) ? this : dev().xferQueue();
@@ -2309,7 +2306,7 @@ void VirtualGPU::submitKernel(amd::NDRangeKernelCommand& vcmd) {
 
     queue->profilingBegin(vcmd);
 
-    static_cast<KernelBlitManager&>(queue->blitMgr()).RunGwsInit(counter);
+    static_cast<KernelBlitManager&>(queue->blitMgr()).RunGwsInit(workgroups);
     queue->addBarrier(RgpSqqtBarrierReason::PostDeviceEnqueue);
 
     // Submit kernel to HW
diff --git a/rocclr/runtime/device/rocm/rocschedcl.cpp b/rocclr/runtime/device/rocm/rocschedcl.cpp
index d0ca5651cd..d699761e32 100644
--- a/rocclr/runtime/device/rocm/rocschedcl.cpp
+++ b/rocclr/runtime/device/rocm/rocschedcl.cpp
@@ -13,23 +13,13 @@ __kernel void scheduler(__global void* params) {
   __amd_scheduler_rocm(params);
 }
 \n);
+
 const char* GwsInitSourceCode = BLIT_KERNEL(
 \n
+extern void __ockl_gws_init(uint nwm1, uint rid);
+\n
 __kernel void gwsInit(uint value) {
-    unsigned int m0_backup, new_m0;
-    __asm__ __volatile__(
-        "s_mov_b32 %0 m0\n"
-        "v_readfirstlane_b32 %1 %2\n"
-        "s_nop 0\n"
-        "s_mov_b32 m0 %1\n"
-        "s_nop 0\n"
-        "ds_gws_init %3 offset:0 gds\n"
-        "s_waitcnt lgkmcnt(0) expcnt(0)\n"
-        "s_mov_b32 m0 %0\n"
-        "s_nop 0"
-        : "=s"(m0_backup), "=s"(new_m0)
-        : "v"(0 << 0x10), "{v0}"(value - 1)
-        : "memory");
+  __ockl_gws_init(value, 0);
 }
 \n);
 
diff --git a/rocclr/runtime/device/rocm/rocvirtual.cpp b/rocclr/runtime/device/rocm/rocvirtual.cpp
index 7b21bfd32c..fda6732fff 100644
--- a/rocclr/runtime/device/rocm/rocvirtual.cpp
+++ b/rocclr/runtime/device/rocm/rocvirtual.cpp
@@ -2203,9 +2203,6 @@ void VirtualGPU::submitKernel(amd::NDRangeKernelCommand& vcmd) {
         workgroups += (vcmd.sizes().global()[i] / vcmd.sizes().local()[i]);
       }
     }
-    uint32_t counter = workgroups *
-      amd::alignUp(vcmd.sizes().local().product(), dev().info().wavefrontWidth_) /
-      dev().info().wavefrontWidth_;
 
     // Get device queue for exclusive GPU access
     VirtualGPU* queue = dev().xferQueue();
@@ -2217,7 +2214,7 @@ void VirtualGPU::submitKernel(amd::NDRangeKernelCommand& vcmd) {
     amd::ScopedLock lock(queue->blitMgr().lockXfer());
     queue->profilingBegin(vcmd);
 
-    static_cast<KernelBlitManager&>(queue->blitMgr()).RunGwsInit(counter);
+    static_cast<KernelBlitManager&>(queue->blitMgr()).RunGwsInit(workgroups);
 
     // Sync AQL packets
     queue->setAqlHeader(kDispatchPacketHeader);