SWDEV-240803 - add hipFuncSetAttribute and hipFuncAttribute

Change-Id: I3f4d67b19d89fd348fa5b884af4a2542ee4aba60


[ROCm/hip commit: 9dfe15a843]

projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md

@@ -103,7 +103,7 @@
 |   **CUDA**                                                |   **HIP**                             |**CUDA version\***|
 |-----------------------------------------------------------|---------------------------------------|:----------------:|
 | `cudaFuncGetAttributes`                                   |`hipFuncGetAttributes`                 |
-| `cudaFuncSetAttribute`                                    |                                       | 9.0              |
+| `cudaFuncSetAttribute`                                    |`hipFuncSetAttribute`                  | 9.0              |
 | `cudaFuncSetCacheConfig`                                  |`hipFuncSetCacheConfig`                |
 | `cudaFuncSetSharedMemConfig`                              |                                       |
 | `cudaGetParameterBuffer`                                  |                                       |

projects/hip/include/hip/hcc_detail/hip_runtime_api.h

@@ -290,6 +290,14 @@ typedef enum hipJitOption {
     hipJitOptionNumOptions
 } hipJitOption;
 
+/**
+ * @warning On AMD devices and some Nvidia devices, these hints and controls are ignored.
+ */
+typedef enum hipFuncAttribute {
+    hipFuncAttributeMaxDynamicSharedMemorySize = 8,
+    hipFuncAttributePreferredSharedMemoryCarveout = 9,
+    hipFuncAttributeMax
+} hipFuncAttribute;
 
 /**
  * @warning On AMD devices and some Nvidia devices, these hints and controls are ignored.
@@ -528,6 +536,21 @@ hipError_t hipDeviceGetCacheConfig(hipFuncCache_t* cacheConfig);
 hipError_t hipDeviceGetLimit(size_t* pValue, enum hipLimit_t limit);
 
 
+/**
+ * @brief Set attribute for a specific function
+ *
+ * @param [in] func;
+ * @param [in] attr;
+ * @param [in] value;
+ *
+ * @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 hipFuncSetAttribute(const void* func, hipFuncAttribute attr, int value);
+
 /**
  * @brief Set Cache configuration for a specific function
  *

projects/hip/include/hip/nvcc_detail/hip_runtime_api.h

@@ -185,6 +185,7 @@ typedef cudaStream_t hipStream_t;
 typedef cudaIpcEventHandle_t hipIpcEventHandle_t;
 typedef cudaIpcMemHandle_t hipIpcMemHandle_t;
 typedef enum cudaLimit hipLimit_t;
+typedef enum cudaFuncAttribute hipFuncAttribute;
 typedef enum cudaFuncCache hipFuncCache_t;
 typedef CUcontext hipCtx_t;
 typedef enum cudaSharedMemConfig hipSharedMemConfig;
@@ -1007,6 +1008,10 @@ inline static hipError_t hipDeviceGetCacheConfig(hipFuncCache_t* pCacheConfig) {
     return hipCUDAErrorTohipError(cudaDeviceGetCacheConfig(pCacheConfig));
 }
 
+inline static hipError_t hipFuncSetAttribute(const void* func, hipFuncAttribute attr, int value) {
+    return hipCUDAErrorTohipError(cudaFuncSetAttribute(func, attr, value));
+}
+
 inline static hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig) {
     return hipCUDAErrorTohipError(cudaDeviceSetCacheConfig(cacheConfig));
 }

projects/hip/rocclr/hip_hcc.def.in

@@ -55,6 +55,7 @@ hipExtModuleLaunchKernel
 hipExtLaunchKernel
 hipFree
 hipFreeArray
+hipFuncSetAttribute
 hipFuncSetCacheConfig
 hipGetDevice
 hipGetDeviceCount

projects/hip/rocclr/hip_hcc.map.in

@@ -56,6 +56,7 @@ global:
     hipExtLaunchKernel;
     hipFree;
     hipFreeArray;
+    hipFuncSetAttribute;
     hipFuncSetCacheConfig;
     hipGetDevice;
     hipGetDeviceCount;

projects/hip/rocclr/hip_module.cpp

@@ -193,6 +193,14 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func)
   HIP_RETURN(hipSuccess);
 }
 
+hipError_t hipFuncSetAttribute ( const void* func, hipFuncAttribute attr, int value ) {
+  HIP_INIT_API(hipFuncSetAttribute, func, attr, value);
+
+  // No way to set function attribute 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

@@ -1408,6 +1408,14 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func)
     return ihipLogStatus(hipSuccess);
 }
 
+hipError_t hipFuncSetAttribute(const void* func, hipFuncAttribute attr, int value) {
+    HIP_INIT_API(hipFuncSetCacheConfig, func, attr, value);
+
+    // 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);