AI/rocm-systems
2
0
Fork 0
Código Issues Pull Requests Ações Pacotes Projetos Versões Wiki Atividade

Added HIP_SYMBOL macro to act as a wrapper between HCC and NVCC symbol name parameters

Ver Código-Fonte 
Change-Id: I008d028b1e29d5a00d0e449af388216396ad2f75


[ROCm/clr commit: 2a55ae10e8]
Esse commit está contido em:
Aditya Atluri
2016-10-13 10:31:56 -05:00
pai a208e9be9c
commit 7b180aff34
3 arquivos alterados com 7 adições e 15 exclusões
Baixar Arquivo de Patch Baixar Arquivo de Diff Expandir todos os arquivos Colapsar todos os arquivos
projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h
+1
Ver Arquivo
@@ -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);
projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h
+3 -3
Ver Arquivo
@@ -45,7 +45,7 @@ kernelName<<<numblocks,numthreads,memperblock,streamId>>>(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<<<numblocks,numthreads,memperblock,streamId>>>(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<<<numblocks,numthreads,memperblock,streamId>>>(0, ##__VA_ARGS__);\
#define HIP_DYNAMIC_SHARED_ATTRIBUTE
#endif
projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp
+3 -12
Ver Arquivo
@@ -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<NUM;i++) {