From d71c0d10de647ad114233f43b4279c21544c74ba Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Tue, 11 Oct 2016 13:50:31 -0500 Subject: [PATCH] changed hipTestDeviceSymbol test to compile for both nvcc and hcc path Change-Id: I041770ad59d4f88d0c8d27d90cdc8a799935ada1 --- include/hip/nvcc_detail/hip_runtime_api.h | 2 +- tests/src/deviceLib/hipTestDeviceSymbol.cpp | 16 +++++++++++++++- 2 files changed, 16 insertions(+), 2 deletions(-) diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index f59585ef08..838465bd0d 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/include/hip/nvcc_detail/hip_runtime_api.h @@ -301,7 +301,7 @@ inline static hipError_t hipMemcpyToSymbol(const void* symbol, const void* src, return hipCUDAErrorTohipError(cudaMemcpyToSymbol(symbol, src, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(copyType))); } -inline static hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t sizeBytes, size_t offset = 0, hipMemcpyKind copyType = hipMemcpyHostToDevice, hipStream_t stream) { +inline static hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t sizeBytes, size_t offset, hipMemcpyKind copyType, hipStream_t stream) { return hipCUDAErrorTohipError(cudaMemcpyToSymbolAsync(symbol, src, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(copyType))); } diff --git a/tests/src/deviceLib/hipTestDeviceSymbol.cpp b/tests/src/deviceLib/hipTestDeviceSymbol.cpp index 6a71295bd2..8317a53990 100644 --- a/tests/src/deviceLib/hipTestDeviceSymbol.cpp +++ b/tests/src/deviceLib/hipTestDeviceSymbol.cpp @@ -24,8 +24,13 @@ THE SOFTWARE. #define NUM 1024 #define SIZE 1024*4 +#ifdef __HIP_PLATFORM_HCC__ __attribute__((address_space(1))) int global[NUM]; +#endif +#ifdef __HIP_PLATFORM_NVCC__ +__device__ int global[NUM]; +#endif __global__ void Assign(hipLaunchParm lp, int* Out) { @@ -47,7 +52,12 @@ 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 hipStreamSynchronize(stream); hipLaunchKernel(Assign, dim3(1,1,1), dim3(NUM,1,1), 0, 0, Ad); hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost); @@ -60,8 +70,12 @@ 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 hipLaunchKernel(Assign, dim3(1,1,1), dim3(NUM,1,1), 0, 0, Ad); hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost); for(unsigned i=0;i