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


[ROCm/clr commit: f23c1a1499]

projects/clr/hipamd/hip_prof_gen.py

@@ -333,7 +333,7 @@ def generate_prof_header(f, api_map, opts_map):
   
   # Generating the callbacks ID enumaration
   f.write('\n// Return HIP API string\n')
-  f.write('static inline const char* hip_api_name(const uint32_t& id) {\n')
+  f.write('static inline const char* hip_api_name(const uint32_t id) {\n')
   f.write('  switch(id) {\n')
   for name in api_map.keys():
     f.write('    case HIP_API_ID_' + name + ': return "' +  name + '";\n')
@@ -353,7 +353,10 @@ def generate_prof_header(f, api_map, opts_map):
     if len(args) != 0:
       f.write('    struct {\n')
       for arg_tuple in args:
-        f.write('      ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n')
+	if arg_tuple[0] == "hipLimit_t":
+	    f.write('      enum ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n')
+	else:
+            f.write('      ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n')
       f.write('    } ' + name + ';\n')
   f.write(
   '  } args;\n' +

projects/clr/hipamd/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" {

projects/clr/hipamd/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);

projects/clr/hipamd/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;
 

projects/clr/hipamd/tests/src/dynamicLoading/complex_loading_behavior.cpp

@@ -21,7 +21,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD_CMD: libfoo_amd %hc %S/%s -o libfoo.so -fPIC -lpthread -shared -DTEST_SHARED_LIBRARY EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD_CMD: libfoo_amd %hc %S/%s -o libfoo.so -Xcompiler -fPIC -lpthread -shared -DTEST_SHARED_LIBRARY EXCLUDE_HIP_PLATFORM nvcc
  * BUILD_CMD: libfoo_nvidia %hc %S/%s -o libfoo.so -Xcompiler -fPIC -lpthread -shared -DTEST_SHARED_LIBRARY EXCLUDE_HIP_PLATFORM hcc
  * BUILD_CMD: %t %hc %S/%s -o %T/%t -ldl
  * TEST: %t

projects/clr/hipamd/tests/src/g++/hipMalloc.cpp

@@ -1,7 +1,34 @@
+/*
+ * Copyright (c) 2019-2020 Advanced Micro Devices, Inc. All rights reserved.
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ * */
+
+/* HIT_START
+ * BUILD_CMD: hipMalloc %cxx -D__HIP_PLATFORM_HCC__ -I%hip-path/include -I/opt/rocm/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/%t -std=c++11 EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t EXCLUDE_HIP_PLATFORM nvcc
+ * HIT_END
+ */
+
+
 #include <hip/hip_runtime_api.h>
 #include <iostream>
 
 int main() {
     int* Ad;
     hipMalloc((void**)&Ad, 1024);
+    std::cout<<"PASSED!"<<std::endl;   
 }

projects/clr/hipamd/tests/src/gcc/hipMalloc.c

@@ -1,3 +1,29 @@
+/*
+ * Copyright (c) 2019-2020 Advanced Micro Devices, Inc. All rights reserved.
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ * */
+
+/* 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
+ * TEST: %t EXCLUDE_HIP_PLATFORM nvcc
+ * HIT_END
+ */
+
+
 #include<hip/hip_runtime_api.h>
 #include<stdio.h>
 
@@ -5,4 +31,5 @@ int main()
 {
   int *Ad;
   hipMalloc((void**)&Ad, 1024);
+  printf("PASSED!\n");
 }