From 8b918b065a8dac81094d6fb7881bd23bc076e048 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Fri, 26 Aug 2016 10:32:01 -0500 Subject: [PATCH] Added NVCC support and name changes - Added NVCC support for module APIs - Changed hipFunction and hipModule data types to hipFunction_t and hipModule_t - Created new intenal ihipModuleGetFunction as it is used twice - Changed test to match with the new data types Change-Id: I300a1c7fd40ed7065b1b8b9de97e3a06b96ed729 --- include/hcc_detail/hip_runtime_api.h | 16 +++++++------- include/nvcc_detail/hip_runtime_api.h | 31 +++++++++++++++++++++++++++ src/hip_module.cpp | 22 ++++++++++++------- tests/src/hipModule.cpp | 7 ++---- 4 files changed, 55 insertions(+), 21 deletions(-) diff --git a/include/hcc_detail/hip_runtime_api.h b/include/hcc_detail/hip_runtime_api.h index 1c50c33359..a2b5587ee6 100644 --- a/include/hcc_detail/hip_runtime_api.h +++ b/include/hcc_detail/hip_runtime_api.h @@ -52,9 +52,9 @@ typedef struct ihipDevice_t *hipDevice_t; typedef struct ihipStream_t *hipStream_t; -typedef struct ihipModule_t *hipModule; +typedef struct ihipModule_t *hipModule_t; -typedef struct ihipFunction_t *hipFunction; +typedef struct ihipFunction_t *hipFunction_t; typedef void* hipDeviceptr; @@ -1134,17 +1134,17 @@ hipError_t hipDeviceGetFromId(hipDevice_t *device, int deviceId); hipError_t hipDriverGetVersion(int *driverVersion) ; -hipError_t hipModuleLoad(hipModule *module, const char *fname); +hipError_t hipModuleLoad(hipModule_t *module, const char *fname); -hipError_t hipModuleUnload(hipModule module); +hipError_t hipModuleUnload(hipModule_t module); -hipError_t hipModuleGetFunction(hipFunction *function, hipModule module, const char *kname); +hipError_t hipModuleGetFunction(hipFunction_t *function, hipModule_t module, const char *kname); -hipError_t hipModuleGetGlobal(hipDeviceptr *dptr, size_t *bytes, hipModule hmod, const char *name); +hipError_t hipModuleGetGlobal(hipDeviceptr *dptr, size_t *bytes, hipModule_t hmod, const char *name); -hipError_t hipModuleLoadData(hipModule *module, const void *image); +hipError_t hipModuleLoadData(hipModule_t *module, const void *image); -hipError_t hipLaunchModuleKernel(hipFunction f, +hipError_t hipLaunchModuleKernel(hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, diff --git a/include/nvcc_detail/hip_runtime_api.h b/include/nvcc_detail/hip_runtime_api.h index 3ef12a294b..d017e37495 100644 --- a/include/nvcc_detail/hip_runtime_api.h +++ b/include/nvcc_detail/hip_runtime_api.h @@ -64,6 +64,10 @@ typedef CUcontext hipCtx_t; typedef CUsharedconfig hipSharedMemConfig; typedef CUfunc_cache hipFuncCache; typedef CUdevice hipDevice_t; +typedef CUModule hipModule_t; +typedef CUFunction hipFunction_t; +typedef CUdeviceptr hipDeviceptr; + //typedef cudaChannelFormatDesc hipChannelFormatDesc; #define hipChannelFormatDesc cudaChannelFormatDesc @@ -581,6 +585,33 @@ inline static hipError_t hipCtxGetFlags ( unsigned int* flags ) return hipCUResultTohipError(cuCtxGetFlags ( flags )); } +inline static hipError_t hipModuleGetFunction(hipFunction_t *function, + hipModule_t module, const char *kname) +{ + return hipCUResultTohipError(cuModuleGetFunction(function, module, kname)); +} + +inline static hipError_t hipModuleGetGlobal(hipDeviceptr *dptr, size_t *bytes, + hipModule_t hmod, const char* name) +{ + return hipCUResultTohipError(cuModuleGetGlobal(dptr, bytes, hmod, name)); +} + +inline static hipError_t hipModuleLoad(hipModule_t *module, const char* fname) +{ + return hipCUResultTohipError(cuModuleLoad(module, fname)); +} + +inline static hipError_t hipModuleLoadData(hipModule_t *module, const void *image) +{ + return hipCUResultTohipError(cuModuleLoadData(module, image)); +} + +inline static hipError_t hipModuleUnload(hipModule_t hmod) +{ + return hipCUResultTohipError(cuModuleUnload(hmod)); +} + #ifdef __cplusplus } #endif diff --git a/src/hip_module.cpp b/src/hip_module.cpp index e98cf7458d..ba9e94d9fb 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -98,7 +98,7 @@ uint64_t ElfSize(const void *emi){ return total_size; } -hipError_t hipModuleLoad(hipModule *module, const char *fname){ +hipError_t hipModuleLoad(hipModule_t *module, const char *fname){ HIP_INIT_API(fname); hipError_t ret = hipSuccess; *module = new ihipModule_t; @@ -159,7 +159,7 @@ hipError_t hipModuleLoad(hipModule *module, const char *fname){ return ret; } -hipError_t hipModuleUnload(hipModule hmod){ +hipError_t hipModuleUnload(hipModule_t hmod){ hipError_t ret = hipSuccess; hsa_status_t status = hsa_executable_destroy(hmod->executable); if(status != HSA_STATUS_SUCCESS){ret = hipErrorInvalidValue; } @@ -169,7 +169,7 @@ hipError_t hipModuleUnload(hipModule hmod){ return ret; } -hipError_t hipModuleGetFunction(hipFunction *func, hipModule hmod, const char *name){ +hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const char *name){ HIP_INIT_API(name); auto ctx = ihipGetTlsDefaultCtx(); hipError_t ret = hipSuccess; @@ -217,7 +217,13 @@ hipError_t hipModuleGetFunction(hipFunction *func, hipModule hmod, const char *n return ret; } -hipError_t hipLaunchModuleKernel(hipFunction f, +hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod, + const char *name) +{ + return ihipModuleGetFunction(hfunc, hmod, name); +} + +hipError_t hipLaunchModuleKernel(hipFunction_t f, uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ, uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ, uint32_t sharedMemBytes, hipStream_t hStream, @@ -283,7 +289,7 @@ Kernel argument preparation. hipError_t hipModuleGetGlobal(hipDeviceptr *dptr, size_t *bytes, - hipModule hmod, const char* name){ + hipModule_t hmod, const char* name){ hipError_t ret = hipSuccess; if(dptr == NULL || bytes == NULL){ return hipErrorInvalidValue; @@ -292,15 +298,15 @@ hipError_t hipModuleGetGlobal(hipDeviceptr *dptr, size_t *bytes, return hipErrorNotInitialized; } else{ - hipFunction func; - hipModuleGetFunction(&func, hmod, name); + hipFunction_t func; + ihipModuleGetFunction(&func, hmod, name); *bytes = PrintSymbolSizes(hmod->ptr, name) + sizeof(amd_kernel_code_t); *dptr = reinterpret_cast(func->kernel); return ret; } } -hipError_t hipModuleLoadData(hipModule *module, const void *image){ +hipError_t hipModuleLoadData(hipModule_t *module, const void *image){ hipError_t ret; if(image == NULL || module == NULL){ return hipErrorNotInitialized; diff --git a/tests/src/hipModule.cpp b/tests/src/hipModule.cpp index 784a01c1d0..94daa2fc5a 100644 --- a/tests/src/hipModule.cpp +++ b/tests/src/hipModule.cpp @@ -21,9 +21,6 @@ THE SOFTWARE. #include #include #include -#include -#include -#include #include #define LEN 64 @@ -53,8 +50,8 @@ int main(){ hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice); hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice); - hipModule Module; - hipFunction Function; + hipModule_t Module; + hipFunction_t Function; hipModuleLoad(&Module, fileName); hipModuleGetFunction(&Function, Module, kernel_name); hipStream_t stream;