From d42822325ff9543ee2eeca4788655cb89fd2e6b0 Mon Sep 17 00:00:00 2001 From: Todd tiantuo Li Date: Tue, 18 Aug 2020 05:19:01 -0700 Subject: [PATCH] SWDEV-240803 - add hipFuncSetSharedMemConfig Change-Id: I160b04677b3e7b99b3981ae7ecc84a0e3811d5e8 [ROCm/hip commit: 99eb486937a063b28d8a07b83ca892bada74685d] --- ...CUDA_Runtime_API_functions_supported_by_HIP.md | 2 +- .../hip/include/hip/hcc_detail/hip_runtime_api.h | 15 ++++++++++++++- .../hip/include/hip/nvcc_detail/hip_runtime_api.h | 4 ++++ projects/hip/rocclr/hip_hcc.def.in | 1 + projects/hip/rocclr/hip_hcc.map.in | 1 + projects/hip/rocclr/hip_module.cpp | 8 ++++++++ projects/hip/src/hip_module.cpp | 8 ++++++++ 7 files changed, 37 insertions(+), 2 deletions(-) diff --git a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md index 20c3aaa9a1..b9c9708ccb 100644 --- a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md @@ -105,7 +105,7 @@ | `cudaFuncGetAttributes` |`hipFuncGetAttributes` | | `cudaFuncSetAttribute` |`hipFuncSetAttribute` | 9.0 | | `cudaFuncSetCacheConfig` |`hipFuncSetCacheConfig` | -| `cudaFuncSetSharedMemConfig` | | +| `cudaFuncSetSharedMemConfig` |`hipFuncSetSharedMemConfig` | | `cudaGetParameterBuffer` | | | `cudaGetParameterBufferV2` | | | `cudaLaunchKernel` |`hipLaunchKernel` | diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index 36b1ddf79e..7f81d73eb0 100755 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -312,7 +312,6 @@ typedef enum hipFuncCache_t { hipFuncCachePreferEqual, ///< prefer equal size L1 cache and shared memory } hipFuncCache_t; - /** * @warning On AMD devices and some Nvidia devices, these hints and controls are ignored. */ @@ -566,6 +565,20 @@ hipError_t hipFuncSetAttribute(const void* func, hipFuncAttribute attr, int valu */ hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t config); +/** + * @brief Set shared memory configuation for a specific function + * + * @param [in] func + * @param [in] config + * + * @returns #hipSuccess, #hipErrorInvalidDeviceFunction, #hipErrorInvalidValue + * + * Note: AMD devices and some Nvidia GPUS do not support shared cache banking, and the hint is + * ignored on those architectures. + * + */ +hipError_t hipFuncSetSharedMemConfig(const void* func, hipSharedMemConfig config); + /** * @brief Returns bank width of shared memory for current device * diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index e7050422cb..faa0bf7d7b 100755 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -1016,6 +1016,10 @@ inline static hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig) { return hipCUDAErrorTohipError(cudaDeviceSetCacheConfig(cacheConfig)); } +inline static hipError_t hipFuncSetSharedMemConfig(const void* func, hipSharedMemConfig config) { + return hipCUDAErrorTohipError(cudaFuncSetSharedMemConfig(func, config)); +} + inline static const char* hipGetErrorString(hipError_t error) { return cudaGetErrorString(hipErrorToCudaError(error)); } diff --git a/projects/hip/rocclr/hip_hcc.def.in b/projects/hip/rocclr/hip_hcc.def.in index 5573e581e3..7e767a1881 100755 --- a/projects/hip/rocclr/hip_hcc.def.in +++ b/projects/hip/rocclr/hip_hcc.def.in @@ -57,6 +57,7 @@ hipFree hipFreeArray hipFuncSetAttribute hipFuncSetCacheConfig +hipFuncSetSharedMemConfig hipGetDevice hipGetDeviceCount hipGetDeviceProperties diff --git a/projects/hip/rocclr/hip_hcc.map.in b/projects/hip/rocclr/hip_hcc.map.in index 62b89a42eb..989cbca40b 100755 --- a/projects/hip/rocclr/hip_hcc.map.in +++ b/projects/hip/rocclr/hip_hcc.map.in @@ -58,6 +58,7 @@ global: hipFreeArray; hipFuncSetAttribute; hipFuncSetCacheConfig; + hipFuncSetSharedMemConfig; hipGetDevice; hipGetDeviceCount; hipGetDeviceProperties; diff --git a/projects/hip/rocclr/hip_module.cpp b/projects/hip/rocclr/hip_module.cpp index 4648a6a3a2..146c6f3829 100755 --- a/projects/hip/rocclr/hip_module.cpp +++ b/projects/hip/rocclr/hip_module.cpp @@ -210,6 +210,14 @@ hipError_t hipFuncSetCacheConfig (const void* func, hipFuncCache_t cacheConfig) HIP_RETURN(hipSuccess); } +hipError_t hipFuncSetSharedMemConfig ( const void* func, hipSharedMemConfig config) { + HIP_INIT_API(hipFuncSetSharedMemConfig, func, config); + + // No way to set Shared Memory config function 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, diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index 0571e0753f..65077d389a 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -1416,6 +1416,14 @@ hipError_t hipFuncSetAttribute(const void* func, hipFuncAttribute attr, int valu return ihipLogStatus(hipSuccess); } +hipError_t hipFuncSetSharedMemConfig(const void* func, hipSharedMemConfig config) { + HIP_INIT_API(hipFuncSetSharedMemConfig, func, config); + + // Nop, AMD does not support setting shared memory size for function. + + return ihipLogStatus(hipSuccess); +} + hipError_t hipFuncGetAttribute(int* value, hipFunction_attribute attrib, hipFunction_t hfunc) { HIP_INIT_API(hipFuncGetAttribute, value, attrib, hfunc);