From c9bf543320ecc4f5d2a65eb49843a665d618fca7 Mon Sep 17 00:00:00 2001 From: foreman Date: Sun, 29 Apr 2018 15:04:45 -0400 Subject: [PATCH] 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: 9f4d0049aa1d22be6b6be3fb3969a71d84f75d30] --- projects/clr/hipamd/api/hip/hip_module.cpp | 7 ++++--- projects/clr/hipamd/api/hip/hip_platform.cpp | 2 +- 2 files changed, 5 insertions(+), 4 deletions(-) diff --git a/projects/clr/hipamd/api/hip/hip_module.cpp b/projects/clr/hipamd/api/hip/hip_module.cpp index 0bcae3551b..268b92284b 100644 --- a/projects/clr/hipamd/api/hip/hip_module.cpp +++ b/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
(extra[1]) + desc.offset_); + kernel->parameters().set(i, desc.size_, reinterpret_cast
(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*/); } } diff --git a/projects/clr/hipamd/api/hip/hip_platform.cpp b/projects/clr/hipamd/api/hip/hip_platform.cpp index 9c9c3b1ecf..7152458d38 100644 --- a/projects/clr/hipamd/api/hip/hip_platform.cpp +++ b/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,