changed hipTestDeviceSymbol test to compile for both nvcc and hcc path
Change-Id: I041770ad59d4f88d0c8d27d90cdc8a799935ada1
[ROCm/clr commit: 8e8939d775]
This commit is contained in:
@@ -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)));
|
||||
}
|
||||
|
||||
|
||||
@@ -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++) {
|
||||
|
||||
مرجع در شماره جدید
Block a user