diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h index efd7771531..01c4128648 100755 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -620,6 +620,7 @@ __device__ static inline void* free(void *ptr) #define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE) #define HIP_KERNEL_NAME(...) __VA_ARGS__ +#define HIP_SYMBOL(X) #X #ifdef __HCC_CPP__ extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_launch_parm *lp, const char *kernelNameStr); diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h index 569d6297bf..c6d8147684 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h @@ -45,7 +45,7 @@ kernelName<<>>(0, ##__VA_ARGS__);\ #define __HIP_ARCH_HAS_GLOBAL_FLOAT_ATOMIC_EXCH__ (__CUDA_ARCH__ >= 110) #define __HIP_ARCH_HAS_SHARED_INT32_ATOMICS__ (__CUDA_ARCH__ >= 120) #define __HIP_ARCH_HAS_SHARED_FLOAT_ATOMIC_EXCH__ (__CUDA_ARCH__ >= 120) -#define __HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__ +#define __HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__ // 64-bit Atomics: #define __HIP_ARCH_HAS_GLOBAL_INT64_ATOMICS__ (__CUDA_ARCH__ >= 200) @@ -92,6 +92,8 @@ kernelName<<>>(0, ##__VA_ARGS__);\ #define hipGridDim_y gridDim.y #define hipGridDim_z gridDim.z +#define HIP_SYMBOL(X) X + /** * extern __shared__ */ @@ -102,5 +104,3 @@ kernelName<<>>(0, ##__VA_ARGS__);\ #define HIP_DYNAMIC_SHARED_ATTRIBUTE #endif - - diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp index 8317a53990..1158bf3f9d 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp @@ -52,12 +52,7 @@ int main() hipStream_t stream; hipStreamCreate(&stream); -#ifdef __HIP_PLATFORM_HCC__ - hipMemcpyToSymbolAsync("global", A, SIZE, 0, hipMemcpyHostToDevice, stream); -#endif -#ifdef __HIP_PLATFORM_NVCC__ - hipMemcpyToSymbolAsync(global, A, SIZE, 0, hipMemcpyHostToDevice, stream); -#endif + hipMemcpyToSymbolAsync(HIP_SYMBOL(global), A, SIZE, 0, hipMemcpyHostToDevice, stream); hipStreamSynchronize(stream); hipLaunchKernel(Assign, dim3(1,1,1), dim3(NUM,1,1), 0, 0, Ad); hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost); @@ -70,12 +65,8 @@ int main() A[i] = -2*i; B[i] = 0; } -#ifdef __HIP_PLATFORM_HCC__ - hipMemcpyToSymbol("global", A, SIZE, 0, hipMemcpyHostToDevice); -#endif -#ifdef __HIP_PLATFORM_NVCC__ - hipMemcpyToSymbol(global, A, SIZE, 0, hipMemcpyHostToDevice); -#endif + + hipMemcpyToSymbol(HIP_SYMBOL(global), A, SIZE, 0, hipMemcpyHostToDevice); hipLaunchKernel(Assign, dim3(1,1,1), dim3(NUM,1,1), 0, 0, Ad); hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost); for(unsigned i=0;i