SWDEV-240803 - add hipFuncSetSharedMemConfig

Change-Id: I160b04677b3e7b99b3981ae7ecc84a0e3811d5e8


[ROCm/hip commit: 99eb486937]

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`                      |

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
  *

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));
 }

projects/hip/rocclr/hip_hcc.def.in

@@ -57,6 +57,7 @@ hipFree
 hipFreeArray
 hipFuncSetAttribute
 hipFuncSetCacheConfig
+hipFuncSetSharedMemConfig
 hipGetDevice
 hipGetDeviceCount
 hipGetDeviceProperties

projects/hip/rocclr/hip_hcc.map.in

@@ -58,6 +58,7 @@ global:
     hipFreeArray;
     hipFuncSetAttribute;
     hipFuncSetCacheConfig;
+    hipFuncSetSharedMemConfig;
     hipGetDevice;
     hipGetDeviceCount;
     hipGetDeviceProperties;

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,

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);