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,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);
 

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

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

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