University:
Boston University
Course:
GRS CS 910 | Directed Study: Systems
Academic year:

2023
Views:

292

Pages:

13
Author:

Daniela Larson

kernels vs. parallel (1)  kernels • More implicit. • Gives the compiler more freedom to find and map parallelism. • Compiler performs parallel analysis and parallelizes what it believes safe.  parallel • More explicit. • Requires analysis by programmer to ensure safe parallelism • Straightforward path from OpenMP kernels vs. parallel (2)  Parallelize a code block with two loops: kernels #pragma acc kernels { for (i=0; i MAX_TEMP_ERROR && iteration <= max_iterations ) { A Loop for Jacobi iterations. 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]); Loops for computing a new solution } dt = 0.0; 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++; } Loops for updating the solution and finding the max error. Laplace Solver (serial Fortran) do while ( dt > max_temp_error .and. iteration <= max_iterations) 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 dt=0.0 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 iteration = iteration+1 enddo A Loop for Jacobi iterations. Loops for computing a new solution Loops for updating the solution and finding the max error. Exercise 2: Laplace Solver in OpenACC  Provided a serial code (in C or Fortran) for solving the two-dimensional Laplace equation, parallelize it using OpenACC directives. Then compare the performance between the serial code and the OpenACC code. • Hints: 1. Find the “hot spots”, the most time-consuming parts of the code. Usually they are loops. 2. Analyze parallelism. Which loops are parallelizable? 3. What directives should be used? Where to insert the directives? Laplace Solver (OpenACC in C , version 1) while ( dt > MAX_TEMP_ERROR && iteration <= max_iterations ) { This loop is not parallelizable due to data dependency. #pragma acc kernels 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]); } These loops are parallelizable. Create a kernel region and ask the compiler to determine parallelism and 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++; } These loops are parallelizable. Create a kernel region and ask the compiler to determine parallelism and data transfer. Laplace Solver (OpenACC in Fortran, version 1) 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 This loop is not parallelizable due to data dependency. These loops are parallelizable. Create a kernel region and ask the compiler to determine parallelism and data transfer. These loops are parallelizable. Create a kernel region and ask the compiler to determine parallelism and data transfer. Analysis of performance (version 1)  Compare the computation time (for 1000*1000 grid): • Serial code: 17.610445 seconds. • OpenACC code (version 1): 48.796347 seconds  The OpenACC code is much slower than the serial code. What went wrong?  We need to further analyze the parallelism and data transfer. time(us): 25,860,945 61: compute region reached 3372 times 63: kernel launched 3372 times grid: [32x250] block: [32x4] device time(us): total=1,006,028 max=312 min=296 avg=298 elapsed time(us): total=1,149,681 max=862 min=337 avg=340 61: data region reached 6744 times 61: data copyin transfers: 3372 device time(us): total=4,570,063 max=1,378 min=1,353 avg=1,355 69: data copyout transfers: 3372 device time(us): total=4,217,959 max=1,987 min=1,248 avg=1,250 72: compute region reached 3372 times 74: kernel launched 3372 times grid: [32x250] block: [32x4] device time(us): total=1,143,160 max=342 min=325 avg=339 elapsed time(us): total=1,300,500 max=1,128 min=373 avg=385 74: reduction kernel launched 3372 times grid: [1] block: [256] device time(us): total=67,550 max=21 min=20 avg=20 elapsed time(us): total=146,840 max=436 min=42 avg=43 72: data region reached 6744 times 72: data copyin transfers: 6744 device time(us): total=9,567,773 max=1,648 min=1,346 avg=1,418 81: data copyout transfers: 3372 device time(us): total=5,176,980 max=1,553 min=1,534 avg=1,535 Profiling (version 1)  export PGI_ACC_TIME=1 to activate profiling, then run again. • There are four times data transfer between host(CPU) memory and GPU memory in every iteration of the outer while loop. • The total time for data transfer is around 23.6 seconds, which is much larger than the computing time around 2.5 seconds! Analysis of data transfer (version 1)  These data transfer happen every iteration of the outer while loop! while ( dt > MAX_TEMP_ERROR && iteration <= max_iterations ) { #pragma acc kernels for(i = 1; i <= ROWS; i++) Copy in: A and A_new is copied from host to GPU. 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]); } Copy out: A and A_new is copied from GPU to host. dt = 0.0; #pragma acc kernels for(i = 1; i <= ROWS; i++) Copy in: A and A_new is copied from host to GPU. 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++; } Copy out: A and A_new is copied from GPU to host. Data clauses  copy (list): Allocates memory on GPU and copies data from host to GPU when entering region and copies data to the host when exiting region.  copyin(list): Allocates memory on GPU and copies data from host to GPU when entering region.  copyout(list): Allocates memory on GPU and copies data to the host when exiting region.  create(list): Allocates memory on GPU but does not copy.  present(list): Data is already present on GPU. • Syntax for C #pragma acc data copy(a[0:size]) copyin(b[0:size]), copyout(c[0:size]) create(d[0:size]) present(d[0:size]) • Syntax for Fortran !$acc acc data copy(a(0:size)) copyin(b(0:size]), copyout(c(0:size)) create(d(0:size)) present(d(0:size)) !$acc end data • If the compiler can determine the size of arrays, it is unnecessary to specify it explicitly.

Introduction to OpenACC Part 2

Related Documents

Recommended Documents

Description

Related Documents

Free up your schedule!

Take 5 seconds to unlock