Introduction to OpenACC Part 3

Laplace Solver (OpenACC in C , version 2) #pragma acc data copy(A), create(A_new) while ( dt > MAX_TEMP_ERROR && iteration <= max_iterations ) { #pragma acc kernels Create a data region here. A is copied in before the while loop starts and is copied out after the while loop ends. A_new is allocated on GPU memory directly and is unnecessary to be copied to the host memory. for(i = 1; i <= ROWS; i++) for(j = 1; j <= COLUMNS; j++) { A_new[i][j] = 0.25 * (A[i+1][j] + A[i-1][j] + A[i][j+1] + A[i][j-1]); } Create a kernel region here to parallelize the for loops, but there is no data transfer. dt = 0.0; #pragma acc kernels for(i = 1; i <= ROWS; i++) for(j = 1; j <= COLUMNS; j++){ dt = fmax( fabs(A_new[i][j]-A[i][j]), dt); A[i][j] = A_new[i][j]; } iteration++; } Create a kernel region here to parallelize the for loops, but there is no data transfer. Laplace Solver (OpenACC in Fortran , version 2) !$acc data copy(A), create(A_new) do while ( dt > max_temp_error .and. iteration <= max_iterations) !$acc kernels do j=1,columns do i=1,rows A_new(i,j)=0.25*(A(i+1,j)+A(i-1,j)+ A(i,j+1)+A(i,j-1) ) enddo enddo !$acc end kernels dt=0.0 !$acc kernels do j=1,columns do i=1,rows dt = max( abs(A_new(i,j) - A(i,j)), dt ) A(i,j) = A_new(i,j) enddo enddo !$acc end kernels iteration = iteration+1 enddo Create a data region here. A is copied in before the while loop starts and is copied out after the while loop ends. A_new is allocated on GPU memory directly and is unnecessary to be copied to the host memory. Create a kernel region here to parallelize the do loops, but there is no data transfer. Create a kernel region here to parallelize the do loops, but there is no data transfer. time(us): 2,374,331 59: data region reached 2 times 59: data copyin transfers: 1 device time(us): total=1,564 max=1,564 min=1,564 avg=1,564 91: data copyout transfers: 1 device time(us): total=1,773 max=1,773 min=1,773 avg=1,773 63: compute region reached 3372 times 65: kernel launched 3372 times grid: [32x250] block: [32x4] device time(us): total=1,005,947 max=313 min=296 avg=298 elapsed time(us): total=1,102,391 max=946 min=324 avg=326 74: compute region reached 3372 times 74: data copyin transfers: 3372 device time(us): total=20,344 max=16 min=6 avg=6 76: kernel launched 3372 times grid: [32x250] block: [32x4] device time(us): total=1,150,552 max=344 min=327 avg=341 elapsed time(us): total=1,235,344 max=856 min=352 avg=366 76: reduction kernel launched 3372 times grid: [1] block: [256] device time(us): total=67,484 max=21 min=19 avg=20 elapsed time(us): total=151,147 max=358 min=43 avg=44 76: data copyout transfers: 3372 device time(us): total=68,104 max=46 min=17 avg=20 Profiling (version 2)  export PGI_ACC_TIME=1 to activate profiling, then run again. • There are only 2 times data movement (of arrays) in total. • There are data movements for the variable dt, but it not an array and thus the transfer processes cost very little time. • The total time for data movement is around 0.09 second, which is much smaller than the computing time (around 2.5 seconds)! Analysis of performance (version 2)  Compare the computation time (for 1000*1000 grid): • Serial code: 17.610445 seconds. • OpenACC code (version 1): 48.796347 seconds • OpenACC code (version 2): 2.592581 seconds  The OpenACC code (version 2) is around 6.8 times faster than the serial code. Cheers!  The speed-up would be even larger if the size of the problem increase.  The maximum size of GPU memory (typically 6 GB or 12 GB) is much smaller than regular CPU memory (e.g. 128 GB on BU SCC). Reduction  As we can see from the profiling results, a reduction kernel is created by the compiler. What is reduction and why is it necessary?  In the previous example, the variable dt can be modified by multiple workers (warps) simultaneously. This is called a data race condition. If data race happened, an incorrect result will be returned.  To avoid data race, a reduction clause is required to protect the concerned variable.  Fortunately, the compiler is smart enough to create a reduction kernel and avoid the data race automatically! dt = 0.0; #pragma acc kernels for(i = 1; i <= ROWS; i++) for(j = 1; j <= COLUMNS; j++){ dt = fmax( fabs(A_new[i][j]-A[i][j]), dt); A[i][j] = A_new[i][j]; } Laplace Solver (OpenACC in C , version 3) #pragma acc data copy(A), create(A_new) Create a data region here. while ( dt > MAX_TEMP_ERROR && iteration <= max_iterations ) { #pragma acc parallel loops for(i = 1; i <= ROWS; i++) Create a parallel region here and parallelize the for loops. for(j = 1; j <= COLUMNS; j++) { A_new[i][j] = 0.25 * (A[i+1][j] + A[i-1][j] + A[i][j+1] + A[i][j-1]); } dt = 0.0; #pragma parallel loops reduction(max:dt) for(i = 1; i <= ROWS; i++) for(j = 1; j <= COLUMNS; j++){ dt = fmax( fabs(A_new[i][j]-A[i][j]), dt); A[i][j] = A_new[i][j]; } iteration++; } Create a parallel region here and parallelize the for loops. Explicitly specify the reduction operator and variable. Laplace Solver (OpenACC in Fortran , version 3) !$acc data copy(A), create(A_new) do while ( dt > max_temp_error .and. iteration <= max_iterations) !$acc parallel loop do j=1,columns do i=1,rows A_new(i,j)=0.25*(A(i+1,j)+A(i-1,j)+ A(i,j+1)+A(i,j-1) ) enddo enddo !$acc end parallel loop dt=0.0 !$acc parallel loop do j=1,columns do i=1,rows dt = max( abs(A_new(i,j) - A(i,j)), dt ) A(i,j) = A_new(i,j) enddo enddo !$acc end parallel loop iteration = iteration+1 enddo Create a data region here. Create a parallel region here and parallelize the do loops. Create a parallel region here and parallelize the do loops. Explicitly specify the reduction operator and variable. Analysis of performance (version 3)  Compare the computation time (for 1000*1000 grid): • Serial code: 17.610445 seconds. • OpenACC code (version 1): 48.796347 seconds • OpenACC code (version 2): 2.592581 seconds • OpenACC code (version 3): 2.259797 seconds  Using parallel directive is a little faster than using kernel directive in this case (mostly due to different task granularities).  It is a good habit to explicitly specify reduction operators and variables. NVIDIA GPU (CUDA) Task Granularity • GPU device -- CUDA grids: Kennels/grids are assigned to a device. • Streaming Multiprocessor (SM) -- CUDA thread blocks: Blocks are assigned to a SM. • CUDA cores -- CUDA threads: Threads are assigned to a core.  Warp: a unit that consists of 32 threads.  Blocks are divided into warps.  The SM executes threads at warp granularity.  The warp size can be changed in the future. OpenACC Task Granularity • Gang --- block • Worker – warp • Vector – thread  Syntax for C #pragma acc kernels loop gang(n) worker(m) vector(k) #pragma acc parallel loop num_gangs(n) num_workers(m) vector_length(k)  Syntax for Fortran !$acc kernels loop gang(n) worker(m) vector(k) !$acc parallel loop num_gangs(n) num_workers(m) vector_length(k) Appendix B: More options for requesting GPU resources on SCC  To request 4 CPU cores and 1 GPU -pe omp 4 -l gpus=0.25  To request 12 CPU cores and 1 GPU (e.g. for budge node) -pe omp 12 -l gpus=0.08  To request a whole budge node (12 CPU cores and 8 GPUs) -pe omp 12 -l gpus=0.6  To request 2 nodes with 12 CPU cores and 8 GPUs on each node -pe mpi_12_tasks_per_node 24 -l gpus=0.6  To request 1 node with 2 K40 GPUs -pe omp 16 –l gpus=0.125 –l gpu_type=K40m What is not covered • Architecture of GPU • Advanced OpenACC (vector, worker, gang, synchronization, etc) • Using OpenACC with CUDA • Using OpenACC with OpenMP (to use a few GPUs on one node) • Using OpenACC with MPI (to use many GPUs on multiple nodes)