AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

changed hipTestDeviceSymbol test to compile for both nvcc and hcc path

Browse Source 
Change-Id: I041770ad59d4f88d0c8d27d90cdc8a799935ada1
This commit is contained in:
Aditya Atluri
2016-10-11 13:50:31 -05:00
parent e3712197ca
commit d71c0d10de
2 changed files with 16 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
tests/src/deviceLib/hipTestDeviceSymbol.cpp
+15 -1
View File
@@ -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<NUM;i++) {