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

[HIP] return an error if blockDim exceeds maxThreadsPerBlock

Browse Source
This commit is contained in:
Aryan Salmanpour
2020-03-10 15:26:53 -04:00
parent 7009901fe0
commit 70ef268add
1 changed files with 4 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/src/hip_module.cpp
+4 -2
View File
@@ -502,7 +502,8 @@ hipError_t ihipLaunchCooperativeKernel(const void* f, dim3 gridDim,
stream = ihipSyncAndResolveStream(stream); stream = ihipSyncAndResolveStream(stream);
if (!stream->getDevice()->_props.cooperativeLaunch) { if (!stream->getDevice()->_props.cooperativeLaunch ||
blockDimX.x * blockDimX.y * blockDimX.z > stream->getDevice()->_props.maxThreadsPerBlock) {
return hipErrorInvalidConfiguration; return hipErrorInvalidConfiguration;
} }
@@ -650,7 +651,8 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL
} }
if (lp.gridDim.x == 0 || lp.gridDim.y == 0 || lp.gridDim.z == 0 || if (lp.gridDim.x == 0 || lp.gridDim.y == 0 || lp.gridDim.z == 0 ||
lp.blockDim.x == 0 || lp.blockDim.y == 0 || lp.blockDim.z == 0){ lp.blockDim.x == 0 || lp.blockDim.y == 0 || lp.blockDim.z == 0 ||
lp.blockDim.x * lp.blockDim.y * lp.blockDim.z > currentDevice->_props.maxThreadsPerBlock){
return hipErrorInvalidConfiguration; return hipErrorInvalidConfiguration;
} }
} }