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


[ROCm/hip commit: 8b918b065a]

projects/hip/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,

projects/hip/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

projects/hip/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<void*>(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;

projects/hip/tests/src/hipModule.cpp

@@ -21,9 +21,6 @@ THE SOFTWARE.
 #include<hip_runtime_api.h>
 #include<iostream>
 #include<fstream>
-#include<hsa/hsa.h>
-#include<hsa/amd_hsa_kernel_code.h>
-#include<hsa/hsa_ven_amd_loader.h>
 #include<vector>
 
 #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;