AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-243961 Added overflow check for globalWorkSize

Browse Source 
Change-Id: I153d99398eb4619baad8c0a52466e0d5d5f898ca


[ROCm/clr commit: 8481524416]
This commit is contained in:
Anusha Godavarthy Surya
2020-07-14 06:41:34 -04:00
committed by agodavar
parent 9fa78dd816
commit 66be9e86d5
2 changed files with 62 additions and 27 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/rocclr/hip_platform.cpp
+12 -2
View File
@@ -629,8 +629,18 @@ hipError_t ihipLaunchKernel(const void* hostFunction,
if ((hip_error != hipSuccess) || (func == nullptr)) {
HIP_RETURN(hipErrorInvalidDeviceFunction);
}
HIP_RETURN(ihipModuleLaunchKernel(func, (gridDim.x * blockDim.x), (gridDim.y * blockDim.y),
(gridDim.z * blockDim.z), blockDim.x, blockDim.y, blockDim.z,
size_t globalWorkSizeX = gridDim.x * blockDim.x;
size_t globalWorkSizeY = gridDim.y * blockDim.y;
size_t globalWorkSizeZ = gridDim.z * blockDim.z;
if (globalWorkSizeX > std::numeric_limits<uint32_t>::max() ||
globalWorkSizeY > std::numeric_limits<uint32_t>::max() ||
globalWorkSizeZ > std::numeric_limits<uint32_t>::max()) {
HIP_RETURN(hipErrorInvalidConfiguration);
}
HIP_RETURN(ihipModuleLaunchKernel(func, static_cast<uint32_t>(globalWorkSizeX),
static_cast<uint32_t>(globalWorkSizeY),
static_cast<uint32_t>(globalWorkSizeZ),
blockDim.x, blockDim.y, blockDim.z,
sharedMemBytes, stream, args, nullptr, startEvent, stopEvent,
flags));
}