Fix gcc build on NVCC path (#1661)

* Fix gcc build on NVCC path

* Fix CI build errors

* [dtest] Fix texture and surface obj2D tests


[ROCm/hip commit: e39d7497ec]

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

@@ -66,24 +66,24 @@ typedef enum hipMemcpyKind {
 #define HIP_LIBRARY_PATCH_LEVEL PATCH_LEVEL
 
 // hipTextureAddressMode
-#define hipTextureAddressMode cudaTextureAddressMode
+typedef enum cudaTextureAddressMode hipTextureAddressMode;
 #define hipAddressModeWrap cudaAddressModeWrap
 #define hipAddressModeClamp cudaAddressModeClamp
 #define hipAddressModeMirror cudaAddressModeMirror
 #define hipAddressModeBorder cudaAddressModeBorder
 
 // hipTextureFilterMode
-#define hipTextureFilterMode cudaTextureFilterMode
+typedef enum cudaTextureFilterMode hipTextureFilterMode;
 #define hipFilterModePoint cudaFilterModePoint
 #define hipFilterModeLinear cudaFilterModeLinear
 
 // hipTextureReadMode
-#define hipTextureReadMode cudaTextureReadMode
+typedef enum cudaTextureReadMode hipTextureReadMode;
 #define hipReadModeElementType cudaReadModeElementType
 #define hipReadModeNormalizedFloat cudaReadModeNormalizedFloat
 
 // hipChannelFormatKind
-#define hipChannelFormatKind            cudaChannelFormatKind
+typedef enum cudaChannelFormatKind hipChannelFormatKind;
 #define hipChannelFormatKindSigned      cudaChannelFormatKindSigned
 #define hipChannelFormatKindUnsigned    cudaChannelFormatKindUnsigned
 #define hipChannelFormatKindFloat       cudaChannelFormatKindFloat
@@ -167,7 +167,7 @@ typedef cudaIpcMemHandle_t hipIpcMemHandle_t;
 typedef enum cudaLimit hipLimit_t;
 typedef enum cudaFuncCache hipFuncCache_t;
 typedef CUcontext hipCtx_t;
-typedef cudaSharedMemConfig hipSharedMemConfig;
+typedef enum cudaSharedMemConfig hipSharedMemConfig;
 typedef CUfunc_cache hipFuncCache;
 typedef CUjit_option hipJitOption;
 typedef CUdevice hipDevice_t;
@@ -177,8 +177,8 @@ typedef CUdeviceptr hipDeviceptr_t;
 typedef struct cudaArray hipArray;
 typedef struct cudaArray* hipArray_t;
 typedef struct cudaArray* hipArray_const_t;
-typedef cudaFuncAttributes hipFuncAttributes;
+typedef struct cudaFuncAttributes hipFuncAttributes;
-typedef CUfunction_attribute hipFunction_attribute;
+#define hipFunction_attribute CUfunction_attribute
 #define hip_Memcpy2D CUDA_MEMCPY2D
 #define hipMemcpy3DParms cudaMemcpy3DParms
 #define hipArrayDefault cudaArrayDefault
@@ -196,8 +196,8 @@ typedef cudaSurfaceObject_t hipSurfaceObject_t;
 #define hipTextureType3D cudaTextureType3D
 #define hipDeviceMapHost cudaDeviceMapHost
 
-#define hipExtent cudaExtent
+typedef struct cudaExtent hipExtent;
-#define hipPitchedPtr cudaPitchedPtr
+typedef struct cudaPitchedPtr hipPitchedPtr;
 #define make_hipExtent make_cudaExtent
 #define make_hipPos make_cudaPos
 #define make_hipPitchedPtr make_cudaPitchedPtr
@@ -205,10 +205,10 @@ typedef cudaSurfaceObject_t hipSurfaceObject_t;
 #define hipStreamDefault cudaStreamDefault
 #define hipStreamNonBlocking cudaStreamNonBlocking
 
-#define hipChannelFormatDesc cudaChannelFormatDesc
+typedef struct cudaChannelFormatDesc hipChannelFormatDesc;
-#define hipResourceDesc cudaResourceDesc
+typedef struct cudaResourceDesc hipResourceDesc;
-#define hipTextureDesc cudaTextureDesc
+typedef struct cudaTextureDesc hipTextureDesc;
-#define hipResourceViewDesc cudaResourceViewDesc
+typedef struct cudaResourceViewDesc hipResourceViewDesc;
 // adding code for hipmemSharedConfig
 #define hipSharedMemBankSizeDefault cudaSharedMemBankSizeDefault
 #define hipSharedMemBankSizeFourByte cudaSharedMemBankSizeFourByte
@@ -394,7 +394,7 @@ inline static enum cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind ki
     }
 }
 
-inline static cudaTextureAddressMode hipTextureAddressModeToCudaTextureAddressMode(
+inline static enum cudaTextureAddressMode hipTextureAddressModeToCudaTextureAddressMode(
     hipTextureAddressMode kind) {
     switch (kind) {
         case hipAddressModeWrap:
@@ -410,7 +410,7 @@ inline static cudaTextureAddressMode hipTextureAddressModeToCudaTextureAddressMo
     }
 }
 
-inline static cudaTextureFilterMode hipTextureFilterModeToCudaTextureFilterMode(
+inline static enum cudaTextureFilterMode hipTextureFilterModeToCudaTextureFilterMode(
     hipTextureFilterMode kind) {
     switch (kind) {
         case hipFilterModePoint:
@@ -422,7 +422,7 @@ inline static cudaTextureFilterMode hipTextureFilterModeToCudaTextureFilterMode(
     }
 }
 
-inline static cudaTextureReadMode hipTextureReadModeToCudaTextureReadMode(hipTextureReadMode kind) {
+inline static enum cudaTextureReadMode hipTextureReadModeToCudaTextureReadMode(hipTextureReadMode kind) {
     switch (kind) {
         case hipReadModeElementType:
             return cudaReadModeElementType;
@@ -433,7 +433,7 @@ inline static cudaTextureReadMode hipTextureReadModeToCudaTextureReadMode(hipTex
     }
 }
 
-inline static cudaChannelFormatKind hipChannelFormatKindToCudaChannelFormatKind(
+inline static enum cudaChannelFormatKind hipChannelFormatKindToCudaChannelFormatKind(
     hipChannelFormatKind kind) {
     switch (kind) {
         case hipChannelFormatKindSigned:
@@ -510,14 +510,14 @@ inline static hipError_t hipMallocManaged(void** ptr, size_t size, unsigned int
     return hipCUDAErrorTohipError(cudaMallocManaged(ptr, size, flags));
 }
 
-inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc,
+inline static hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
                                         size_t width, size_t height,
                                         unsigned int flags __dparm(hipArrayDefault)) {
     return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags));
 }
 
-inline static hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc* desc,
+inline static hipError_t hipMalloc3DArray(hipArray** array, const hipChannelFormatDesc* desc,
-                            struct hipExtent extent, unsigned int flags) {
+                             hipExtent extent, unsigned int flags) {
     return hipCUDAErrorTohipError(cudaMalloc3DArray(array, desc, extent, flags));
 }
 
@@ -1048,7 +1048,6 @@ inline static hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attribut
     return err;
 }
 
-
 inline static hipError_t hipMemGetInfo(size_t* free, size_t* total) {
     return hipCUDAErrorTohipError(cudaMemGetInfo(free, total));
 }
@@ -1073,7 +1072,6 @@ inline static hipError_t hipEventDestroy(hipEvent_t event) {
     return hipCUDAErrorTohipError(cudaEventDestroy(event));
 }
 
-
 inline static hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags) {
     return hipCUDAErrorTohipError(cudaStreamCreateWithFlags(stream, flags));
 }
@@ -1360,79 +1358,15 @@ inline static hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gri
                                                 kernelParams, extra));
 }
 
-
 inline static hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t cacheConfig) {
     return hipCUDAErrorTohipError(cudaFuncSetCacheConfig(func, cacheConfig));
 }
 
-#ifdef __cplusplus
+inline static hipError_t hipBindTexture(size_t* offset, struct textureReference* tex, const void* devPtr,
-}
+                                        const hipChannelFormatDesc* desc, size_t size __dparm(UINT_MAX)){
-#endif
-
-#ifdef __CUDACC__
-
-template <class T>
-inline static hipError_t hipOccupancyMaxPotentialBlockSize(int* minGridSize, int* blockSize, T func,
-                                                           size_t dynamicSMemSize = 0,
-                                                           int blockSizeLimit = 0) {
-    cudaError_t cerror;
-    cerror = cudaOccupancyMaxPotentialBlockSize(minGridSize, blockSize, func, dynamicSMemSize, blockSizeLimit);
-    return hipCUDAErrorTohipError(cerror);
-}
-
-template <class T, int dim, enum hipTextureReadMode readMode>
-inline static hipError_t hipBindTexture(size_t* offset, const struct texture<T, dim, readMode>& tex,
-                                        const void* devPtr, size_t size = UINT_MAX) {
-    return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, size));
-}
-
-template <class T, int dim, enum hipTextureReadMode readMode>
-inline static hipError_t hipBindTexture(size_t* offset, struct texture<T, dim, readMode>& tex,
-                                        const void* devPtr, const struct hipChannelFormatDesc& desc,
-                                        size_t size = UINT_MAX) {
     return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size));
 }
 
-template <class T, int dim, enum hipTextureReadMode readMode>
-inline static hipError_t hipUnbindTexture(struct texture<T, dim, readMode>* tex) {
-    return hipCUDAErrorTohipError(cudaUnbindTexture(tex));
-}
-
-template <class T, int dim, enum hipTextureReadMode readMode>
-inline static hipError_t hipUnbindTexture(struct texture<T, dim, readMode> &tex) {
-    return hipCUDAErrorTohipError(cudaUnbindTexture(tex));
-}
-
-inline static hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* devPtr,
-                                        const hipChannelFormatDesc* desc, size_t size = UINT_MAX){
-    return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size));
-}
-
-template <class T, int dim, enum hipTextureReadMode readMode>
-inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& tex,
-                                               hipArray_const_t array,
-                                               const struct hipChannelFormatDesc& desc) {
-    return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc));
-}
-
-template <class T, int dim, enum hipTextureReadMode readMode>
-inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode> *tex,
-                                               hipArray_const_t array,
-                                               const struct hipChannelFormatDesc* desc) {
-    return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc));
-}
-
-template <class T, int dim, enum hipTextureReadMode readMode>
-inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& tex,
-                                               hipArray_const_t array) {
-    return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array));
-}
-
-template <class T>
-inline static hipChannelFormatDesc hipCreateChannelDesc() {
-    return cudaCreateChannelDesc<T>();
-}
-
 inline static hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w,
                                                         hipChannelFormatKind f) {
     return cudaCreateChannelDesc(x, y, z, w, hipChannelFormatKindToCudaChannelFormatKind(f));
@@ -1464,7 +1398,7 @@ inline static hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDe
     return hipCUDAErrorTohipError(cudaGetTextureObjectResourceDesc( pResDesc, textureObject));
 }
 
-inline static hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference* texref)
+inline static hipError_t hipGetTextureAlignmentOffset(size_t* offset, const struct textureReference* texref)
 {
     return hipCUDAErrorTohipError(cudaGetTextureAlignmentOffset(offset,texref));
 }
@@ -1473,6 +1407,70 @@ inline static hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_
 {
     return hipCUDAErrorTohipError(cudaGetChannelDesc(desc,array));
 }
+
+
+#ifdef __cplusplus
+}
+#endif
+
+#ifdef __CUDACC__
+
+template <class T>
+inline static hipError_t hipOccupancyMaxPotentialBlockSize(int* minGridSize, int* blockSize, T func,
+                                                           size_t dynamicSMemSize = 0,
+                                                           int blockSizeLimit = 0) {
+    cudaError_t cerror;
+    cerror = cudaOccupancyMaxPotentialBlockSize(minGridSize, blockSize, func, dynamicSMemSize, blockSizeLimit);
+    return hipCUDAErrorTohipError(cerror);
+}
+
+template <class T, int dim, enum cudaTextureReadMode readMode>
+inline static hipError_t hipBindTexture(size_t* offset, const struct texture<T, dim, readMode>& tex,
+                                        const void* devPtr, size_t size = UINT_MAX) {
+    return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, size));
+}
+
+template <class T, int dim, enum cudaTextureReadMode readMode>
+inline static hipError_t hipBindTexture(size_t* offset, struct texture<T, dim, readMode>& tex,
+                                        const void* devPtr, const hipChannelFormatDesc& desc,
+                                        size_t size = UINT_MAX) {
+    return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size));
+}
+
+template <class T, int dim, enum cudaTextureReadMode readMode>
+inline static hipError_t hipUnbindTexture(struct texture<T, dim, readMode>* tex) {
+    return hipCUDAErrorTohipError(cudaUnbindTexture(tex));
+}
+
+template <class T, int dim, enum cudaTextureReadMode readMode>
+inline static hipError_t hipUnbindTexture(struct texture<T, dim, readMode> &tex) {
+    return hipCUDAErrorTohipError(cudaUnbindTexture(tex));
+}
+
+template <class T, int dim, enum cudaTextureReadMode readMode>
+inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& tex,
+                                               hipArray_const_t array,
+                                               const hipChannelFormatDesc& desc) {
+    return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc));
+}
+
+template <class T, int dim, enum cudaTextureReadMode readMode>
+inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode> *tex,
+                                               hipArray_const_t array,
+                                               const hipChannelFormatDesc* desc) {
+    return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc));
+}
+
+template <class T, int dim, enum cudaTextureReadMode readMode>
+inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& tex,
+                                               hipArray_const_t array) {
+    return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array));
+}
+
+template <class T>
+inline static hipChannelFormatDesc hipCreateChannelDesc() {
+    return cudaCreateChannelDesc<T>();
+}
 #endif  //__CUDACC__
 
 #endif  // HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_API_H

projects/hip/tests/src/gcc/hipMalloc.c

@@ -18,8 +18,10 @@
  * */
 
 /* HIT_START
- * BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_HCC__ -I%hip-path/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/%t EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_NVCC__ -I%hip-path/include -I/usr/local/cuda/include %S/%s -o %T/hipMalloc_nv -L/usr/local/cuda/lib64 -lcudart EXCLUDE_HIP_PLATFORM hcc
- * TEST: %t EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_HCC__ -I%hip-path/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/hipMalloc_hcc EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: hipMalloc_nv EXCLUDE_HIP_PLATFORM hcc
+ * TEST: hipMalloc_hcc EXCLUDE_HIP_PLATFORM nvcc
  * HIT_END
  */
 

projects/hip/tests/src/surface/hipSurfaceObj2D.cpp

@@ -53,7 +53,7 @@ int runTest(int argc, char** argv) {
 
     hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice);
 
-    struct hipResourceDesc resDesc;
+    hipResourceDesc resDesc;
     memset(&resDesc, 0, sizeof(resDesc));
     resDesc.resType = hipResourceTypeArray;
     resDesc.res.array.array = hipArray;
@@ -62,7 +62,7 @@ int runTest(int argc, char** argv) {
     hipCreateSurfaceObject(&surfaceObject, &resDesc);
 
     hipMallocArray(&hipOutArray, &channelDesc, width, height);
-    struct hipResourceDesc resOutDesc;
+    hipResourceDesc resOutDesc;
     memset(&resOutDesc, 0, sizeof(resOutDesc));
     resOutDesc.resType = hipResourceTypeArray;
     resOutDesc.res.array.array = hipOutArray;

projects/hip/tests/src/texture/hipTextureObj2D.cpp

@@ -53,13 +53,13 @@ int runTest(int argc, char** argv) {
 
     hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice);
 
-    struct hipResourceDesc resDesc;
+    hipResourceDesc resDesc;
     memset(&resDesc, 0, sizeof(resDesc));
     resDesc.resType = hipResourceTypeArray;
     resDesc.res.array.array = hipArray;
 
     // Specify texture object parameters
-    struct hipTextureDesc texDesc;
+    hipTextureDesc texDesc;
     memset(&texDesc, 0, sizeof(texDesc));
     texDesc.addressMode[0] = hipAddressModeWrap;
     texDesc.addressMode[1] = hipAddressModeWrap;