P4 to Git Change 1547669 by lmoriche@lmoriche_opencl_dev2 on 2018/04/29 14:53:31

SWDEV-145570 - [HIP] Fix kernel disptach for HCC compiled programs.

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#5 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#8 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rockernel.cpp#34 edit


[ROCm/clr commit: 9f4d0049aa]

projects/clr/hipamd/api/hip/hip_module.cpp

@@ -151,7 +151,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
   }
 
   size_t globalWorkOffset[3] = {0};
-  size_t globalWorkSize[3] = { gridDimX, gridDimY, gridDimZ };
+  size_t globalWorkSize[3] = { gridDimX * blockDimX, gridDimY * blockDimY, gridDimZ * blockDimZ};
   size_t localWorkSize[3] = { blockDimX, blockDimY, blockDimZ };
   amd::NDRangeContainer ndrange(3, globalWorkOffset, globalWorkSize, localWorkSize);
   amd::Command::EventWaitList waitList;
@@ -161,10 +161,11 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
     const amd::KernelParameterDescriptor& desc = signature.at(i);
     if (kernelParams == nullptr) {
       assert(extra);
-      kernel->parameters().set(i, desc.size_, reinterpret_cast<address>(extra[1]) + desc.offset_);
+      kernel->parameters().set(i, desc.size_, reinterpret_cast<address>(extra[1]) + desc.offset_,
+                               desc.type_ == T_POINTER/*svmBound*/);
     } else {
       assert(!extra);
-      kernel->parameters().set(i, desc.size_, kernelParams[i]);
+      kernel->parameters().set(i, desc.size_, kernelParams[i], desc.type_ == T_POINTER/*svmBound*/);
     }
   }
 

projects/clr/hipamd/api/hip/hip_platform.cpp

@@ -417,7 +417,7 @@ void hipLaunchKernelGGLImpl(
 
   const auto it = functions().find(function_address);
   if (it == functions().cend())
-    return;
+    assert(0);
 
   hipModuleLaunchKernel(it->second,
     numBlocks.x, numBlocks.y, numBlocks.z,