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/hip commit: 9ea54f7797]
This commit is contained in:
Anusha Godavarthy Surya
2020-07-14 06:41:34 -04:00
committed by agodavar
parent 2c90a43675
commit 2cd41daf44
2 changed files with 62 additions and 27 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/rocclr/hip_module.cpp
+50 -25
View File
@@ -193,16 +193,17 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func)
HIP_RETURN(hipSuccess);
}
hipError_t ihipModuleLaunchKernel(hipFunction_t f,
uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ,
hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ,
uint32_t sharedMemBytes, hipStream_t hStream,
void **kernelParams, void **extra,
hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags = 0,
uint32_t params = 0, uint32_t gridId = 0, uint32_t numGrids = 0,
uint64_t prevGridSum = 0, uint64_t allGridSum = 0, uint32_t firstDevice = 0) {
HIP_INIT_API(ihipModuleLaunchKernel, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags, params);
HIP_INIT_API(ihipModuleLaunchKernel, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ,
blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra, startEvent,
stopEvent, flags, params);
hip::DeviceFunc* function = hip::DeviceFunc::asFunction(f);
amd::Kernel* kernel = function->kernel();
@@ -229,7 +230,8 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f,
int block_size = blockDimX * blockDimY * blockDimZ;
hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor(
&num_blocks, &max_blocks_per_grid, &best_block_size, device, f, block_size, sharedMemBytes, true);
if (((gridDimX * gridDimY * gridDimZ) / block_size) > unsigned(max_blocks_per_grid)) {
if (((globalWorkSizeX * globalWorkSizeY * globalWorkSizeZ) / block_size) >
unsigned(max_blocks_per_grid)) {
return hipErrorCooperativeLaunchTooLarge;
}
}
@@ -243,7 +245,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f,
}
size_t globalWorkOffset[3] = {0};
size_t globalWorkSize[3] = { gridDimX, gridDimY, gridDimZ };
size_t globalWorkSize[3] = { globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ };
size_t localWorkSize[3] = { blockDimX, blockDimY, blockDimZ };
amd::NDRangeContainer ndrange(3, globalWorkOffset, globalWorkSize, localWorkSize);
amd::Command::EventWaitList waitList;
@@ -290,11 +292,11 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f,
command->enqueue();
if(startEvent != nullptr) {
if (startEvent != nullptr) {
eStart->addMarker(queue, command, false);
command->retain();
}
if(stopEvent != nullptr) {
if (stopEvent != nullptr) {
eStop->addMarker(queue, command, false);
command->retain();
}
@@ -313,8 +315,17 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream,
kernelParams, extra);
HIP_RETURN(ihipModuleLaunchKernel(f, gridDimX * blockDimX, gridDimY * blockDimY, gridDimZ * blockDimZ,
size_t globalWorkSizeX = gridDimX * blockDimX;
size_t globalWorkSizeY = gridDimY * blockDimY;
size_t globalWorkSizeZ = gridDimZ * blockDimZ;
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(f, static_cast<uint32_t>(globalWorkSizeX),
static_cast<uint32_t>(globalWorkSizeY),
static_cast<uint32_t>(globalWorkSizeZ),
blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream, kernelParams, extra, nullptr, nullptr));
}
@@ -337,37 +348,37 @@ hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t gridDimX,
uint32_t gridDimY, uint32_t gridDimZ,
hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
uint32_t blockDimX, uint32_t blockDimY,
uint32_t blockDimZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams, void** extra,
hipEvent_t startEvent,
hipEvent_t stopEvent)
{
HIP_INIT_API(hipHccModuleLaunchKernel, f, gridDimX, gridDimY, gridDimZ,
HIP_INIT_API(hipHccModuleLaunchKernel, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ,
blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream,
kernelParams, extra, startEvent, stopEvent);
HIP_RETURN(ihipModuleLaunchKernel(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ,
HIP_RETURN(ihipModuleLaunchKernel(f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent));
}
hipError_t hipModuleLaunchKernelExt(hipFunction_t f, uint32_t gridDimX,
uint32_t gridDimY, uint32_t gridDimZ,
hipError_t hipModuleLaunchKernelExt(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
uint32_t blockDimX, uint32_t blockDimY,
uint32_t blockDimZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams, void** extra,
hipEvent_t startEvent,
hipEvent_t stopEvent)
{
HIP_INIT_API(hipModuleLaunchKernelExt, f, gridDimX, gridDimY, gridDimZ,
HIP_INIT_API(hipModuleLaunchKernelExt, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ,
blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream,
kernelParams, extra, startEvent, stopEvent);
HIP_RETURN(ihipModuleLaunchKernel(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ,
HIP_RETURN(ihipModuleLaunchKernel(f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent));
}
@@ -406,8 +417,17 @@ hipError_t hipLaunchCooperativeKernel(const void* f,
int deviceId = ihipGetDevice();
hipFunction_t func = nullptr;
HIP_RETURN_ONFAIL(PlatformState::instance().getStatFunc(&func, f, deviceId));
HIP_RETURN(ihipModuleLaunchKernel(func, gridDim.x * blockDim.x, gridDim.y * blockDim.y, gridDim.z * 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, hStream, kernelParams, nullptr, nullptr, nullptr, 0,
amd::NDRangeKernelCommand::CooperativeGroups));
@@ -481,11 +501,16 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL
result = hipErrorInvalidDeviceFunction;
HIP_RETURN(result);
}
result = ihipModuleLaunchKernel(func,
launch.gridDim.x * launch.blockDim.x,
launch.gridDim.y * launch.blockDim.y,
launch.gridDim.z * launch.blockDim.z,
size_t globalWorkSizeX = launch.gridDim.x * launch.blockDim.x;
size_t globalWorkSizeY = launch.gridDim.y * launch.blockDim.y;
size_t globalWorkSizeZ = launch.gridDim.z * launch.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);
}
result = ihipModuleLaunchKernel(func, static_cast<uint32_t>(globalWorkSizeX),
static_cast<uint32_t>(globalWorkSizeY), static_cast<uint32_t>(globalWorkSizeZ),
launch.blockDim.x, launch.blockDim.y, launch.blockDim.z,
launch.sharedMem, launch.stream, launch.args, nullptr, nullptr, nullptr,
flags, extFlags, i, numDevices, prevGridSize, allGridSize, firstDevice);
projects/hip/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));
}