Merge pull request #364 from lmoriche/master

Add initial support for clang

include/hip/hcc_detail/hip_runtime.h

@@ -525,6 +525,56 @@ do {\
  */
 
 
+#elif defined(__clang__) && defined(__HIP__)
+
+#define hipConfigureCall cudaConfigureCall
+#define hipSetupArgument cudaSetupArgument
+#define hipLaunch cudaLaunch
+
+typedef int hipLaunchParm ;
+
+#define hipLaunchKernel(kernelName, numblocks, numthreads, memperblock, streamId, ...) \
+do {\
+  kernelName<<<numblocks,numthreads,memperblock,streamId>>>(0, ##__VA_ARGS__);\
+} while(0)
+
+#define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \
+do {\
+  kernelName<<<numblocks,numthreads,memperblock,streamId>>>(__VA_ARGS__);\
+} while(0)
+
+#include <hip/hip_runtime_api.h>
+
+#if defined(__cplusplus)
+extern "C" {
+#endif /*__cplusplus*/
+
+hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, hipStream_t stream);
+hipError_t hipSetupArgument(const void *arg, size_t size, size_t offset);
+hipError_t hipLaunch(const void *func);
+
+#if defined(__cplusplus)
+}
+#endif /*__cplusplus*/
+
+#include <__clang_cuda_builtin_vars.h>
+
+#define hipThreadIdx_x threadIdx.x
+#define hipThreadIdx_y threadIdx.y
+#define hipThreadIdx_z threadIdx.z
+
+#define hipBlockIdx_x  blockIdx.x
+#define hipBlockIdx_y  blockIdx.y
+#define hipBlockIdx_z  blockIdx.z
+
+#define hipBlockDim_x  blockDim.x
+#define hipBlockDim_y  blockDim.y
+#define hipBlockDim_z  blockDim.z
+
+#define hipGridDim_x  gridDim.x
+#define hipGridDim_y  gridDim.y
+#define hipGridDim_z  gridDim.z
+
 #endif
 
 #endif//HIP_HCC_DETAIL_RUNTIME_H

include/hip/hcc_detail/host_defines.h

@@ -64,7 +64,19 @@ THE SOFTWARE.
 #define __shared__     tile_static
 #define __constant__   __attribute__((hc))
 
+#elif defined(__clang__) && defined(__HIP__)
+
+#define __host__      __attribute__((host))
+#define __device__    __attribute__((device))
+#define __global__    __attribute__((global))
+#define __shared__    __attribute__((shared))
+#define __constant__  __attribute__((constant))
+
+#define __noinline__     __attribute__((noinline))
+#define __forceinline__  __attribute__((always_inline))
+
 #else
+
 // Non-HCC compiler
 /**
  * Function and kernel markers

include/hip/hip_common.h

@@ -26,12 +26,12 @@ THE SOFTWARE.
 // Common code included at start of every hip file.
 // Auto enable __HIP_PLATFORM_HCC__ if compiling with HCC
 // Other compiler (GCC,ICC,etc) need to set one of these macros explicitly
-#if defined(__HCC__)
+#if defined(__HCC__) || (defined(__clang__) && defined(__HIP__))
 #define __HIP_PLATFORM_HCC__
 #endif //__HCC__
 
 // Auto enable __HIP_PLATFORM_NVCC__ if compiling with NVCC
-#if defined(__NVCC__)
+#if defined(__NVCC__) || (defined(__clang__) && defined(__CUDA__))
 #define __HIP_PLATFORM_NVCC__
 #ifdef __CUDACC__
 #define __HIPCC__