From 1340b3f07fcb5041e18d878ddd734aa0a49b26d5 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Tue, 21 Jul 2020 06:31:13 -0400 Subject: [PATCH 1/3] Adding Anyorder flag to HIP Change-Id: Ie20931541b3febe23fa9ac36ebc0c90de75a5f0a --- include/hip/hcc_detail/hip_runtime_api.h | 3 +++ rocclr/hip_module.cpp | 5 +++++ 2 files changed, 8 insertions(+) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index d4ac271e6c..36b1ddf79e 100755 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -232,6 +232,9 @@ enum hipLimit_t { #define hipCpuDeviceId ((int)-1) #define hipInvalidDeviceId ((int)-2) +// Flags that can be used with hipExtLaunch Set of APIs +#define hipExtAnyOrderLaunch 0x01 ///< AnyOrderLaunch of kernels + /* * @brief HIP Memory Advise values * @enum diff --git a/rocclr/hip_module.cpp b/rocclr/hip_module.cpp index 5c75deb776..051578dc0a 100755 --- a/rocclr/hip_module.cpp +++ b/rocclr/hip_module.cpp @@ -288,6 +288,11 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, profileNDRange = (startEvent != nullptr && stopEvent != nullptr); + // Flag set to 1 signifies that kernel can be launched in anyorder + if (flags & hipExtAnyOrderLaunch) { + params |= amd::NDRangeKernelCommand::AnyOrderLaunch; + } + amd::NDRangeKernelCommand* command = new amd::NDRangeKernelCommand( *queue, waitList, *kernel, ndrange, sharedMemBytes, params, gridId, numGrids, prevGridSum, allGridSum, firstDevice, profileNDRange); From 4f400bc5e90fb60f11f2a558aa678bb8f96d0454 Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Mon, 17 Aug 2020 12:52:40 -0400 Subject: [PATCH 2/3] SWDEV-248057 - fix the calculation of allGridSize used in multi_grid_group() API Change-Id: Ib470094e28dcacaa4769dc5c7ab08924f5b7fa41 --- rocclr/hip_module.cpp | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/rocclr/hip_module.cpp b/rocclr/hip_module.cpp index 051578dc0a..0ddeaedded 100755 --- a/rocclr/hip_module.cpp +++ b/rocclr/hip_module.cpp @@ -461,14 +461,17 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL hipError_t result = hipErrorUnknown; uint64_t allGridSize = 0; + uint32_t blockDims = 0; std::vector mgpu_list(numDevices); for (int i = 0; i < numDevices; ++i) { const hipLaunchParams& launch = launchParamsList[i]; - allGridSize += launch.gridDim.x * launch.gridDim.y * launch.gridDim.z; + blockDims = launch.blockDim.x * launch.blockDim.y * launch.blockDim.z; + allGridSize += launch.gridDim.x * launch.gridDim.y * launch.gridDim.z * + blockDims; // Make sure block dimensions are valid - if (0 == launch.blockDim.x * launch.blockDim.y * launch.blockDim.z) { + if (0 == blockDims) { return hipErrorInvalidConfiguration; } if (launch.stream != nullptr) { From efcb882e1970f694ae95179945d8d79978e3c9f1 Mon Sep 17 00:00:00 2001 From: Todd tiantuo Li Date: Tue, 4 Aug 2020 04:11:35 -0700 Subject: [PATCH 3/3] move hipFuncSetCacheConfig from hip_device to hip_module Change-Id: If5a930e1210e76531f5ba9a0e5f5ec98ad473a19 --- rocclr/hip_device.cpp | 9 --------- rocclr/hip_module.cpp | 9 +++++++++ 2 files changed, 9 insertions(+), 9 deletions(-) diff --git a/rocclr/hip_device.cpp b/rocclr/hip_device.cpp index b116819846..70548d5328 100644 --- a/rocclr/hip_device.cpp +++ b/rocclr/hip_device.cpp @@ -49,15 +49,6 @@ hipError_t hipDeviceGet(hipDevice_t *device, int deviceId) { HIP_RETURN(hipSuccess); }; -hipError_t hipFuncSetCacheConfig (const void* func, hipFuncCache_t cacheConfig) { - - HIP_INIT_API(hipFuncSetCacheConfig, cacheConfig); - - // No way to set cache config yet. - - HIP_RETURN(hipSuccess); -} - hipError_t hipDeviceTotalMem (size_t *bytes, hipDevice_t device) { HIP_INIT_API(hipDeviceTotalMem, bytes, device); diff --git a/rocclr/hip_module.cpp b/rocclr/hip_module.cpp index 0ddeaedded..4648a6a3a2 100755 --- a/rocclr/hip_module.cpp +++ b/rocclr/hip_module.cpp @@ -201,6 +201,15 @@ hipError_t hipFuncSetAttribute ( const void* func, hipFuncAttribute attr, int va HIP_RETURN(hipSuccess); } +hipError_t hipFuncSetCacheConfig (const void* func, hipFuncCache_t cacheConfig) { + + HIP_INIT_API(hipFuncSetCacheConfig, cacheConfig); + + // No way to set cache config yet. + + HIP_RETURN(hipSuccess); +} + hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ,