From f8b0c964628f86bfd9db6004eebe974db1e169c8 Mon Sep 17 00:00:00 2001 From: Laurent Morichetti Date: Wed, 28 Feb 2018 12:31:26 -0800 Subject: [PATCH 1/3] Add initial support for clang --- include/hip/hcc_detail/hip_runtime.h | 52 ++++++++++++++++++++++++++- include/hip/hcc_detail/host_defines.h | 12 +++++++ include/hip/hip_common.h | 4 +-- 3 files changed, 65 insertions(+), 3 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 924e774af0..cff8888140 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -41,7 +41,57 @@ THE SOFTWARE. #include #endif//__cplusplus -#if __HCC__ +#if defined(__clang__) && defined(__HIPCC__) + +#define hipConfigureCall cudaConfigureCall +#define hipSetupArgument cudaSetupArgument +#define hipLaunch cudaLaunch + +typedef int hipLaunchParm ; + +#define hipLaunchKernel(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ +do {\ + kernelName<<>>(0, ##__VA_ARGS__);\ +} while(0) + +#define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ +do {\ + kernelName<<>>(__VA_ARGS__);\ +} while(0) + +#include + +#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 + +#elif __HCC__ // Define NVCC_COMPAT for CUDA compatibility #define NVCC_COMPAT diff --git a/include/hip/hcc_detail/host_defines.h b/include/hip/hcc_detail/host_defines.h index a7acdfccf7..30ce36afb1 100644 --- a/include/hip/hcc_detail/host_defines.h +++ b/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(__HIPCC__) + +#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 diff --git a/include/hip/hip_common.h b/include/hip/hip_common.h index da8ec4a55d..58c10cd83e 100644 --- a/include/hip/hip_common.h +++ b/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(__HIPCC__)) #define __HIP_PLATFORM_HCC__ #endif //__HCC__ // Auto enable __HIP_PLATFORM_NVCC__ if compiling with NVCC -#if defined(__NVCC__) +#if defined(__NVCC__) || (defined(__clang__) && defined(__CUDACC__)) #define __HIP_PLATFORM_NVCC__ #ifdef __CUDACC__ #define __HIPCC__ From 2bc7adfa62d48bdfd44b4236df2e882b9a6353e9 Mon Sep 17 00:00:00 2001 From: Laurent Morichetti Date: Wed, 28 Feb 2018 14:12:09 -0800 Subject: [PATCH 2/3] Fixes a build error with hcc --- include/hip/hcc_detail/hip_runtime.h | 102 +++++++++++++-------------- 1 file changed, 51 insertions(+), 51 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index cff8888140..2201f8aae0 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -41,57 +41,7 @@ THE SOFTWARE. #include #endif//__cplusplus -#if defined(__clang__) && defined(__HIPCC__) - -#define hipConfigureCall cudaConfigureCall -#define hipSetupArgument cudaSetupArgument -#define hipLaunch cudaLaunch - -typedef int hipLaunchParm ; - -#define hipLaunchKernel(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ -do {\ - kernelName<<>>(0, ##__VA_ARGS__);\ -} while(0) - -#define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ -do {\ - kernelName<<>>(__VA_ARGS__);\ -} while(0) - -#include - -#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 - -#elif __HCC__ +#if __HCC__ // Define NVCC_COMPAT for CUDA compatibility #define NVCC_COMPAT @@ -575,6 +525,56 @@ do {\ */ +#elif defined(__clang__) && defined(__HIPCC__) + +#define hipConfigureCall cudaConfigureCall +#define hipSetupArgument cudaSetupArgument +#define hipLaunch cudaLaunch + +typedef int hipLaunchParm ; + +#define hipLaunchKernel(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ +do {\ + kernelName<<>>(0, ##__VA_ARGS__);\ +} while(0) + +#define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ +do {\ + kernelName<<>>(__VA_ARGS__);\ +} while(0) + +#include + +#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 From 65361649143a3d41a6d5a4cc006160bdc9eea4cc Mon Sep 17 00:00:00 2001 From: Laurent Morichetti Date: Wed, 28 Feb 2018 14:20:55 -0800 Subject: [PATCH 3/3] Use __HIP__ instead of __HIPCC__, and __CUDA__ instead of __CUDACC__ --- include/hip/hcc_detail/hip_runtime.h | 2 +- include/hip/hcc_detail/host_defines.h | 2 +- include/hip/hip_common.h | 4 ++-- 3 files changed, 4 insertions(+), 4 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 2201f8aae0..58f741188d 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -525,7 +525,7 @@ do {\ */ -#elif defined(__clang__) && defined(__HIPCC__) +#elif defined(__clang__) && defined(__HIP__) #define hipConfigureCall cudaConfigureCall #define hipSetupArgument cudaSetupArgument diff --git a/include/hip/hcc_detail/host_defines.h b/include/hip/hcc_detail/host_defines.h index 30ce36afb1..479ce07264 100644 --- a/include/hip/hcc_detail/host_defines.h +++ b/include/hip/hcc_detail/host_defines.h @@ -64,7 +64,7 @@ THE SOFTWARE. #define __shared__ tile_static #define __constant__ __attribute__((hc)) -#elif defined(__clang__) && defined(__HIPCC__) +#elif defined(__clang__) && defined(__HIP__) #define __host__ __attribute__((host)) #define __device__ __attribute__((device)) diff --git a/include/hip/hip_common.h b/include/hip/hip_common.h index 58c10cd83e..fafac1d9ba 100644 --- a/include/hip/hip_common.h +++ b/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__) || (defined(__clang__) && defined(__HIPCC__)) +#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__) || (defined(__clang__) && defined(__CUDACC__)) +#if defined(__NVCC__) || (defined(__clang__) && defined(__CUDA__)) #define __HIP_PLATFORM_NVCC__ #ifdef __CUDACC__ #define __HIPCC__