基础知识
host: traditional CPU
devices: GPU processors with a massive unmber of arithmetic units
main source of performance: data parallelism
very fine grain parallelism enabled by hardware thread scheduling!
cuda threads take very few clock cycles to generate and schedule, due to efficient hardware support.
Single instruction, multiple thread (SIMT) -> low overhead threading
Single program, multiple data (SPMD)
GPU execute blocks in a SIMD-like approach. 32 operations from the same thread block are started in parallel. If not a multiple of 32, then some d=functional units be idle.
API
cudaMalloc()
cudaFree()
cudaMemcpy()
blockDim, threadIdx, blockIdx
function<<<dimGrid, dimBlock>>>
dim3
dimensions of thread blocks: multiples of 32 (due to hardware efficiency reasons)
__syncthreads() : for threads in the same block
_ _ shared _ _
workshop
dim3 NumBlock = {(M % 64 == 0)? 64 : 1 + M / 64, 1, 1};
__syncthreads() ;
while (mod <= blockDim . x ) {
if ( threadIdx.x % mod == 0) {
sum[ threadIdx.x] += sum[ threadIdx.x + mod / 2 ] ;
}
mod <<= 1 ;
__syncthreads();
}
for ( int i = 0 ; i < ceil(N/BLOCK SIZE ) ; i++) {
for ( j = 0 ; j < ceil(BLOCK_SIZE / blockDim.x); j++) {
if ( i * BLOCK_SIZE + threadIdx.x < N) && threadIdx.x + j * blockDim . x < BLOCK_SIZE)
vec_shared [threadIdx.x + j * blockDim.x ] = vec [ i * BLOCK_SIZE + threadIdx.x + j* blockDim.x ] ;
}
__syncthreads() ;
if ( tid < M) {
for ( int j = 0 ; j < BLOCK_SIZE ; j++){
sum += mat [ tid * N + i * BLOCK_SIZE + j ] *vec_shared[j] ;
}
}
__syncthreads() ;
}
cuda streams
A stream in CUDA is a sequence of operations that execute on the device in the order in which they are issued by the host code.
The default stream is different from other streams because it is a synchronizing stream with respect to operations on the device:
- no operation in the default stream will begin until all previously issued operations in any stream on the device have completed
- an operation in the default stream must complete before any other operation (in any stream on the device) will begin.
The asynchronous behavior of kernel launches (from host’s perspective) makes overlapping device and host computation very simple. Whether the host function or device kernel completes first does not affect the subsequent device-to-host transfer, which will begin only after the kernel completes.
cudaStream_t stream1;
cudaError_t result;
result = cudaStreamCreate(&stream1) ;
cudaMemcpyAsync(d_a, a, N, cudaMemcpyHostToDevice, stream1);
increment<<<1, N, 0, stream1>>>(d_a);
cudaStreamQuery(stream1)
cudaDeviceSynchronize();
result = cudaStreamDestroy(stream1) ;