Enabled gcc for hip host code (#1214)

* Enabled gcc for hip host code

* Adding tests for hip code + (gcc & g++), without kernels

* Excluding nvcc platforms for gcc and g++ tests + Addressing review comments

* minor code clean-up

* Add rocm include path

* Added relative path for library

* Hiding non supported functions for gcc

* Incorporating review comments

include/hip/hcc_detail/hip_runtime_api.h

@@ -45,7 +45,7 @@ THE SOFTWARE.
 #include <hip/hcc_detail/hip_texture_types.h>
 #include <hip/hcc_detail/hip_surface_types.h>
 
-#if !__HIP_VDI__
+#if !__HIP_VDI__ && defined(__cplusplus)
 #include <hsa/hsa.h>
 #include <hip/hcc_detail/program_state.hpp>
 #endif
@@ -79,9 +79,11 @@ THE SOFTWARE.
   #define __dparm(x)
 #endif
 
+#ifdef __cplusplus
 namespace hip_impl {
 hipError_t hip_init();
 }  // namespace hip_impl
+#endif
 
 // Structure definitions:
 #ifdef __cplusplus
@@ -1459,6 +1461,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName,
 #else
 hipError_t hipModuleGetGlobal(void**, size_t*, hipModule_t, const char*);
 
+#ifdef __cplusplus //Start : Not supported in gcc 
 namespace hip_impl {
 inline
 __attribute__((visibility("hidden")))
@@ -1466,6 +1469,7 @@ hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
                                           const char* name);
 } // Namespace hip_impl.
 
+
 /**
  *  @brief Copies the memory address of symbol @p symbolName to @p devPtr
  *
@@ -1504,15 +1508,18 @@ hipError_t hipGetSymbolSize(size_t* size, const void* symbolName) {
     void* devPtr = nullptr;
     return hip_impl::read_agent_global_from_process(&devPtr, size, (const char*)symbolName);
 }
+#endif // End : Not supported in gcc
 
 #if defined(__cplusplus)
 } // extern "C"
 #endif
 
+#ifdef __cplusplus
 namespace hip_impl {
 hipError_t hipMemcpyToSymbol(void*, const void*, size_t, size_t, hipMemcpyKind,
                              const char*);
 } // Namespace hip_impl.
+#endif
 
 #if defined(__cplusplus)
 extern "C" {
@@ -1541,6 +1548,7 @@ extern "C" {
  * hipMemcpyFromArrayAsync, hipMemcpy2DFromArrayAsync, hipMemcpyToSymbolAsync,
  * hipMemcpyFromSymbolAsync
  */
+#ifdef __cplusplus
 inline
 __attribute__((visibility("hidden")))
 hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src,
@@ -1554,11 +1562,13 @@ hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src,
     return hip_impl::hipMemcpyToSymbol(dst, src, sizeBytes, offset, kind,
                                        (const char*)symbolName);
 }
+#endif
 
 #if defined(__cplusplus)
 } // extern "C"
 #endif
 
+#ifdef __cplusplus
 namespace hip_impl {
 hipError_t hipMemcpyToSymbolAsync(void*, const void*, size_t, size_t,
                                   hipMemcpyKind, hipStream_t, const char*);
@@ -1567,6 +1577,7 @@ hipError_t hipMemcpyFromSymbol(void*, const void*, size_t, size_t,
 hipError_t hipMemcpyFromSymbolAsync(void*, const void*, size_t, size_t,
                                     hipMemcpyKind, hipStream_t, const char*);
 } // Namespace hip_impl.
+#endif
 
 #if defined(__cplusplus)
 extern "C" {
@@ -1597,6 +1608,8 @@ extern "C" {
  * hipMemcpyFromArrayAsync, hipMemcpy2DFromArrayAsync, hipMemcpyToSymbolAsync,
  * hipMemcpyFromSymbolAsync
  */
+
+#ifdef __cplusplus //Start : Not supported in gcc
 inline
 __attribute__((visibility("hidden")))
 hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src,
@@ -1641,6 +1654,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName,
                                               stream,
                                               (const char*)symbolName);
 }
+#endif // End : Not supported in gcc
 
 #endif // __HIP_VDI__
 /**
@@ -2604,6 +2618,7 @@ hipError_t hipFuncGetAttributes(struct hipFuncAttributes* attr, const void* func
 } // extern "C"
 #endif
 
+#ifdef __cplusplus
 namespace hip_impl {
     class agent_globals_impl;
     class agent_globals {
@@ -2635,6 +2650,7 @@ namespace hip_impl {
         return get_agent_globals().read_agent_global_from_process(dptr, bytes, name);
     }
 } // Namespace hip_impl.
+#endif
 
 #if defined(__cplusplus)
 extern "C" {

include/hip/hcc_detail/hip_vector_types.h

@@ -849,6 +849,7 @@ __MAKE_VECTOR_TYPE__(longlong, long long);
 __MAKE_VECTOR_TYPE__(float, float);
 __MAKE_VECTOR_TYPE__(double, double);
 
+#ifdef __cplusplus
 #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \
     static inline __device__ __host__ \
     type make_##type(comp x) { type r{x}; return r; }
@@ -867,6 +868,26 @@ __MAKE_VECTOR_TYPE__(double, double);
         type r{x, y, z, w}; \
         return r; \
     }
+#else
+ #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \
+     static inline __device__ __host__ \
+     type make_##type(comp x) { type r; r.x =x; return r; }
+
+ #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \
+     static inline __device__ __host__ \
+     type make_##type(comp x, comp y) { type r; r.x=x; r.y=y; return r; }
+
+ #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \
+     static inline __device__ __host__ \
+     type make_##type(comp x, comp y, comp z) { type r; r.x=x; r.y=y; r.z=z; return r; }
+
+ #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \
+     static inline __device__ __host__ \
+     type make_##type(comp x, comp y, comp z, comp w) { \
+         type r; r.x=x; r.y=y; r.z=z; r.w=w; \
+         return r; \
+     }
+#endif
 
 DECLOP_MAKE_ONE_COMPONENT(unsigned char, uchar1);
 DECLOP_MAKE_TWO_COMPONENT(unsigned char, uchar2);

include/hip/hip_runtime_api.h

@@ -255,6 +255,8 @@ typedef enum __HIP_NODISCARD hipError_t {
         1071,    ///< Produced when the IPC memory attach failed from ROCr.
     hipErrorAssert =
         1081,    ///< Produced when the kernel calls assert.
+    hipErrorNotSupported = 
+        1082,    ///< Produced when the hip API is not supported/implemented 
     hipErrorTbd  ///< Marker that more error codes are needed.
 } hipError_t;