Add support for hipFuncGetAttribute (#1279)

* Add support for hipFunGetAttribute

* Support NVCC path

* Test using sample module_api_global

* Try fixing CI build failure due to hip_prof_gen scan

* Fix for CI build issue

* Resolve conflict

* Rebase and resolve conflicts with master

* Fix build error

* Fix NVCC path build error

include/hip/hcc_detail/driver_types.h

@@ -319,4 +319,18 @@ static inline struct hipExtent make_hipExtent(size_t w, size_t h, size_t d) {
     return e;
 }
 
+typedef enum hipFunction_attribute {
+    HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK,
+    HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES,
+    HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES,
+    HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES,
+    HIP_FUNC_ATTRIBUTE_NUM_REGS,
+    HIP_FUNC_ATTRIBUTE_PTX_VERSION,
+    HIP_FUNC_ATTRIBUTE_BINARY_VERSION,
+    HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA,
+    HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES,
+    HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT,
+    HIP_FUNC_ATTRIBUTE_MAX
+}hipFunction_attribute;
+
 #endif

include/hip/hcc_detail/hip_runtime_api.h

@@ -2603,16 +2603,27 @@ hipError_t hipModuleUnload(hipModule_t module);
 hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, const char* kname);
 
 /**
- * @bried Find out attributes for a given function.
+ * @brief Find out attributes for a given function.
  *
  * @param [out] attr
  * @param [in] func
  *
- * @returns hipSuccess, hipErrorInvalidDeviceFunction
+ * @returns hipSuccess, hipErrorInvalidValue, hipErrorInvalidDeviceFunction
  */
 
 hipError_t hipFuncGetAttributes(struct hipFuncAttributes* attr, const void* func);
 
+/**
+ * @brief Find out a specific attribute for a given function.
+ *
+ * @param [out] value
+ * @param [in]  attrib
+ * @param [in]  hfunc
+ *
+ * @returns hipSuccess, hipErrorInvalidValue, hipErrorInvalidDeviceFunction
+ */
+hipError_t hipFuncGetAttribute(int* value, hipFunction_attribute attrib, hipFunction_t hfunc);
+
 #if !__HIP_VDI__
 #if defined(__cplusplus)
 } // extern "C"

include/hip/nvcc_detail/hip_runtime_api.h

@@ -162,6 +162,7 @@ typedef CUdeviceptr hipDeviceptr_t;
 typedef struct cudaArray hipArray;
 typedef struct cudaArray* hipArray_const_t;
 typedef cudaFuncAttributes hipFuncAttributes;
+typedef CUfunction_attribute hipFunction_attribute;
 #define hip_Memcpy2D CUDA_MEMCPY2D
 #define hipMemcpy3DParms cudaMemcpy3DParms
 #define hipArrayDefault cudaArrayDefault
@@ -197,6 +198,19 @@ typedef cudaSurfaceObject_t hipSurfaceObject_t;
 #define hipSharedMemBankSizeFourByte cudaSharedMemBankSizeFourByte
 #define hipSharedMemBankSizeEightByte cudaSharedMemBankSizeEightByte
 
+//Function Attributes
+#define HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK
+#define HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES
+#define HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES CU_FUNC_ATTRIBUTE_CONST_SIZE_BYTES
+#define HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES
+#define HIP_FUNC_ATTRIBUTE_NUM_REGS CU_FUNC_ATTRIBUTE_NUM_REGS
+#define HIP_FUNC_ATTRIBUTE_PTX_VERSION CU_FUNC_ATTRIBUTE_PTX_VERSION
+#define HIP_FUNC_ATTRIBUTE_BINARY_VERSION CU_FUNC_ATTRIBUTE_BINARY_VERSION
+#define HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA CU_FUNC_ATTRIBUTE_CACHE_MODE_CA
+#define HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES
+#define HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT CU_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT
+#define HIP_FUNC_ATTRIBUTE_MAX CU_FUNC_ATTRIBUTE_MAX
+
 inline static hipError_t hipCUDAErrorTohipError(cudaError_t cuError) {
     switch (cuError) {
         case cudaSuccess:
@@ -1197,6 +1211,10 @@ inline static hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const voi
     return hipCUDAErrorTohipError(cudaFuncGetAttributes(attr, func));
 }
 
+inline static hipError_t hipFuncGetAttribute (int* value, hipFunction_attribute attrib, hipFunction_t hfunc) {
+    return hipCUResultTohipError(cuFuncGetAttribute(value, attrib, hfunc));
+}
+
 inline static hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod,
                                             const char* name) {
     return hipCUResultTohipError(cuModuleGetGlobal(dptr, bytes, hmod, name));

samples/0_Intro/module_api_global/runKernel.cpp

@@ -124,6 +124,11 @@ int main() {
     {
         hipFunction_t Function;
         HIP_CHECK(hipModuleGetFunction(&Function, Module, "test_globals"));
+        int val =-1;
+        HIP_CHECK(hipFuncGetAttribute(&val, HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES,Function));
+        printf("Shared Size Bytes = %d\n",val);
+        HIP_CHECK(hipFuncGetAttribute(&val, HIP_FUNC_ATTRIBUTE_NUM_REGS, Function));
+        printf("Num Regs = %d\n",val);
         HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config));
 
         hipMemcpyDtoH(B, Bd, SIZE);

src/hip_module.cpp

@@ -803,6 +803,52 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func)
     return ihipLogStatus(hipSuccess);
 }
 
+hipError_t hipFuncGetAttribute(int* value, hipFunction_attribute attrib, hipFunction_t hfunc)
+{
+    HIP_INIT_API(hipFuncGetAttribute, value, attrib, hfunc);
+    using namespace hip_impl;
+    
+    hipError_t retVal = hipSuccess;
+    if (!value) return ihipLogStatus(hipErrorInvalidValue);
+    hipFuncAttributes attr{};
+    attr = make_function_attributes(tls, *hfunc);
+    switch(attrib) {
+        case HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES:
+            *value = (int) attr.sharedSizeBytes;
+            break;
+        case HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK:
+            *value = attr.maxThreadsPerBlock;
+            break;
+        case HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES:
+            *value = (int) attr.constSizeBytes;
+            break;
+        case HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES:
+            *value = (int) attr.localSizeBytes;
+            break;
+        case HIP_FUNC_ATTRIBUTE_NUM_REGS:
+            *value = attr.numRegs;
+            break;
+        case HIP_FUNC_ATTRIBUTE_PTX_VERSION:
+            *value = attr.ptxVersion;
+            break;
+        case HIP_FUNC_ATTRIBUTE_BINARY_VERSION:
+            *value = attr.binaryVersion;
+            break;
+        case HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA:
+            *value = attr.cacheModeCA;
+            break;
+        case HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES:
+            *value = attr.maxDynamicSharedSizeBytes;
+            break;
+        case HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT:
+            *value = attr.preferredShmemCarveout;
+            break;
+        default:
+             retVal = hipErrorInvalidValue;
+    }
+    return ihipLogStatus(retVal);
+}
+
 hipError_t ihipModuleLoadData(TlsData *tls, hipModule_t* module, const void* image) {
     using namespace hip_impl;