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__