From 5c03a593e770e0f8456e5cd01de596bfc943775d Mon Sep 17 00:00:00 2001 From: Saleel Kudchadker Date: Tue, 15 Oct 2024 03:57:01 +0000 Subject: [PATCH] SWDEV-491375 - Optimize multithreaded dispatches - Fix typo Change-Id: If4c68455dcfa03fee18cb4720e8b5b438642703c [ROCm/clr commit: 0f2342bc133ecc7506de2a0d57c44ab082df9ca0] --- projects/clr/hipamd/src/hip_module.cpp | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index 7e670be022..9dfb745737 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -203,7 +203,7 @@ hipError_t hipFuncSetAttribute(const void* func, hipFuncAttribute attr, int valu *(hip::getCurrentDevice()->devices()[0]))); if (attr == hipFuncAttributeMaxDynamicSharedMemorySize) { - if ((value < 0) || (value > (d_kernel->workGroupInfo()->availableLDSSize_ - + if ((value < 0) || (value > (d_kernel->workGroupInfo()->availableLDSSize_ - d_kernel->workGroupInfo()->localMemSize_))) { HIP_RETURN(hipErrorInvalidValue); } @@ -420,8 +420,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, } hip::DeviceFunc* function = hip::DeviceFunc::asFunction(f); amd::Kernel* kernel = function->kernel(); - - amd::ScopedLock lock (DEBUG_HIP_KERNARG_COPY_OPT ? &function->dflock_ : nullptr); + amd::ScopedLock lock (DEBUG_HIP_KERNARG_COPY_OPT ? nullptr : &function->dflock_); hipError_t status = ihipLaunchKernel_validate( f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, blockDimX, blockDimY, blockDimZ, @@ -734,7 +733,7 @@ hipError_t hipLaunchKernel_common(const void* hostFunction, dim3 gridDim, dim3 b hipError_t hipLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDim, void** args, size_t sharedMemBytes, hipStream_t stream) { HIP_INIT_API(hipLaunchKernel, hostFunction, gridDim, blockDim, args, sharedMemBytes, stream); - HIP_RETURN(hipLaunchKernel_common(hostFunction, gridDim, blockDim, args, sharedMemBytes, stream)); + HIP_RETURN_DURATION(hipLaunchKernel_common(hostFunction, gridDim, blockDim, args, sharedMemBytes, stream)); } hipError_t hipLaunchKernel_spt(const void* hostFunction, dim3 gridDim, dim3 blockDim,