From 6ce86f409d3635b2a5808f74f6809e3192fdee14 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 8 Aug 2019 01:27:41 -0700 Subject: [PATCH] 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 | 14 ++++++ include/hip/hcc_detail/hip_runtime_api.h | 15 +++++- include/hip/nvcc_detail/hip_runtime_api.h | 18 ++++++++ .../0_Intro/module_api_global/runKernel.cpp | 5 ++ src/hip_module.cpp | 46 +++++++++++++++++++ 5 files changed, 96 insertions(+), 2 deletions(-) diff --git a/include/hip/hcc_detail/driver_types.h b/include/hip/hcc_detail/driver_types.h index 18f344d642..510d3d058e 100644 --- a/include/hip/hcc_detail/driver_types.h +++ b/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 diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index e8871912fa..e71ac8d8fb 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/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" diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index fb4987feee..a77d5560d8 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/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)); diff --git a/samples/0_Intro/module_api_global/runKernel.cpp b/samples/0_Intro/module_api_global/runKernel.cpp index 7031b9f898..7aa7f7ba2d 100644 --- a/samples/0_Intro/module_api_global/runKernel.cpp +++ b/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); diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 04cdedbb75..93751167ee 100644 --- a/src/hip_module.cpp +++ b/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;