1 / 28

A Source-to- Source OpenACC compiler for CUDA

A Source-to- Source OpenACC compiler for CUDA . Akihiro Tabuchi †1 Masahiro Nakao †2 Mitsuhisa Sato †1 †1. Graduate School of Systems and Information Engineering, University of Tsukuba †2. Center for Computational Sciences, University of Tsukuba. Outline. Background OpenACC

diella
Download Presentation

A Source-to- Source OpenACC compiler for CUDA

An Image/Link below is provided (as is) to download presentation Download Policy: Content on the Website is provided to you AS IS for your information and personal use and may not be sold / licensed / shared on other websites without getting consent from its author. Content is provided to you AS IS for your information and personal use only. Download presentation by click this link. While downloading, if for some reason you are not able to download a presentation, the publisher may have deleted the file from their server. During download, if you can't get a presentation, the file might be deleted by the publisher.

E N D

Presentation Transcript

  1. A Source-to-Source OpenACC compiler for CUDA Akihiro Tabuchi†1 Masahiro Nakao†2 MitsuhisaSato†1 †1. Graduate School of Systems and Information Engineering, University of Tsukuba †2. Center for Computational Sciences, University of Tsukuba

  2. Outline • Background • OpenACC • Compiler Implementation • Performance Evaluation • Conclusion & Future Work

  3. Background • Accelerator programming model • CUDA (for NVIDIA GPU) • OpenCL (for various accelerators) • Accelerator programming is complex • memory management, kernel function, … • low productivity & low portability • OpenACC is proposed to solve these problems

  4. OpenACC • The directive-based programming model for accelerators • support C, C++ and Fortran • Offloading model • offload a part of code to an accelerator • High productivity • only adding directives • High portability • run on any accelerators as long as the compiler supports it

  5. Example of OpenACC intmain(){ inti; int a[N], b[N], c[N]; /* initialize array ‘a’ and ‘b’ */ #pragma accparallel loop copyin(a,b) copyout(c) for(i = 0; i < N; i++){ c[i] = a[i] + b[i]; } } This directive specifies data transfers and loop offloading and parallelization

  6. Purpose of Research • Designing and implementing an open source OpenACC compiler • Target language : C • Target accelerator : NVIDIA GPU • Source-to-source approach • C + OpenACC → C + CUDA API • This approach enables to leave detailed machine-specific code optimization to the mature CUDA compiler by NVIDIA • The result of compilation is a executable file

  7. Related Work • Commercial compiler • PGI Accelerator compiler • CAPS HMPP • Cray compiler • Open source compiler • accULL • developed at University of La Laguna in Spain • Source-to-source translation • Backend is CUDA and OpenCL • Output is codes and a Makefile

  8. OpenACC directives • parallel • kernels • loop • data • host_data • update • wait • cache • declare • parallel loop • kernels loop (OpenACC specification 1.0)

  9. dataconstruct host memory device memory int a[4];#pragma acc data copy(a) { /* some codes using ‘a’ */ } computation ondevice Data management on Accelerator If an array is specified in “copy” clause … • Device memory allocation • Data transfer from host to device • Data transfer from device to host • Device memory release at the beginning of region at the end of region

  10. Translation of data construct int a[4];#pragma acc data copy(a) { /* some codes using ‘a’ */ } • host address • device address • size • …. int a[4];{ void *_ACC_DEVICE_ADDR_a,*_ACC_HOST_DESC_a; _ACC_gpu_init_data(&_ACC_HOST_DESC_a, &_ACC_DEVICE_ADDR_a, a, 4*sizeof(int)); _ACC_gpu_copy_data(_ACC_HOST_DESC_a, 400); { /* some codes using ‘a’ */ }_ACC_gpu_copy_data(_ACC_HOST_DESC_a, 401);_ACC_gpu_finalize_data(_ACC_HOST_DESC_b); } allocate ‘a’ on GPU copy ‘a’ to GPU from host copy ‘a’ to host from GPU free ‘a’ on GPU

  11. parallel construct #pragma acc parallel num_gangs(1) vector_length(128) { /* codes in parallel region */ } • Codes in parallel region are executed on device • Three levels of parallelism • gang • worker • vector • The number of gang or worker or vector length can be specified by clauses

  12. Translation of parallelconstruct #pragma acc parallel num_gangs(1) vector_length(128) { /* codes in parallel region */ } __global__ static void _ACC_GPU_FUNC_0_DEVICE( ... ) { /* codes in parallel region */ } extern "C” void_ACC_GPU_FUNC_0( … ) { dim3 _ACC_block(1, 1, 1), _ACC_thread(128, 1, 1); _ACC_GPU_FUNC_0_DEVICE<<<_ACC_block,_ACC_thread>>>( ... ); _ACC_GPU_M_BARRIER_KERNEL(); } GPU kernel function kernel launch function

  13. loopconstruct /* inside parallel region */ #pragma acc loop vector for(i = 0; i < 256; i++){ a[i]++; } • Loop construct describes parallelism of loop • Distribute loop iteration among gang, worker or vector • Two or more parallelisms can be specified for a loop • Loops with no loop directive in parallel region is basically executed serially.

  14. Translation of loop construct (1/3) /* inner parallel region */ #pragma acc loop vector for(i = 0; i < N; i++){ a[i]++; } The virtual index is divided and distributed among blocks and/or threads A virtual index which is the same length as loop iteration is prepared Each thread calculates the value of loop variable from the virtual index and executes loop body

  15. Translation of loop construct (2/3) /* inner parallel region */ #pragma acc loop vector for(i = 0; i < N; i++){ a[i]++; } virtual index：_ACC_idx virtual index range : _ACC_init, cond, step calculate the range of virtual index /* inner gpukernel code */ inti, _ACC_idx; int _ACC_init, _ACC_cond, _ACC_step; _ACC_gpu_init_thread_x_iter(&_ACC_init, &_ACC_cond, &_ACC_step, 0, N, 1); for(_ACC_idx = _ACC_init; _ACC_idx < _ACC_cond; _ACC_idx += _ACC_step){ _ACC_gpu_calc_idx(_ACC_idx, &i, 0, N, 1); a[i]++; } virtual index range variables calculate ‘i’ from virtual index loop body

  16. Translation of loop construct(3/3) • Our compiler supports 2D blocking for nested loops • Nested loops are distributed among the 2D blocks in the 2D grid in CUDA (default block size is 16x16) • But it’s not allowed in OpenACC 2.0 and “tile” clause is provided instead 2D Grid #pragma acc loop gang vector for(i = 0; i < N; i++) #pragma acc loop gang vector for(j = 0; j < N; j++) /* … */ distribute 2D Block

  17. Compiler Implementation • Our compiler translates C with OpenACC directives to C with CUDA API • read C code with directives and output translated code • using Omni compiler infrastructure • Omni compiler infrastructure • aset of programs for a source-to-source compiler with code analysis and transformation • supports C and Fortran95

  18. Flow of Compilation Omni Frontend OpenACC translator sample.c sample.xml XcodeML C with OpenACC directives Omni compiler infrastructure acc runtime C compiler sample _tmp.c sample_tmp.o a.out C with ACC API nvcc sample.cu sample.gpu.o CUDA

  19. Performance Evaluation • Benchmark • Matrix multiplication • N-body problem • NAS Parallel Benchmarks – CG • Evaluation environment • 1 node of Cray XK6m-200 • CPU: AMD Opteron Processor 6272 (2.1GHz) • GPU :NVIDIA X2090 (MatMul, N-body) : NVIDIA K20 (NPB CG)

  20. Performance Comparison • Cray compiler • Our compiler • Hand written CUDA • The code is written in CUDA and compiled by NVCC • The code doesn’t use shared memory of GPU • Our compiler (2D-blocking) • The code uses 2D blocking and is compiled by our compiler • This is applied to only matrix multiplication

  21. Matrix multiplication 5.5x 4.6x 1.4x 1.5x The performance of our compiler using 2D-blocking and hand-written CUDA are slightly lower

  22. Matrix multiplication • Our compiler achieves better performance than that of Cray compiler • The PTX code directly generated by Cray compiler has more operations in the innermost loop • Our compiler outputs CUDA code, and NVCC generates more optimized PTX code • 2D-blocking is lower performance • default 2D block size (16x16) is not adequate to this program • the best block size was 512x2 • Hand-written CUDA code also uses 16x16 block

  23. N-body 31x 1.2x 5.4x 0.95x At the small problem size, the performance of our compiler is lower than that of Cray compiler

  24. N-body • At small problem size, the performance became worse • Decline in the utilization of Streaming Multiprocessors(SMs) • A kernel is executed by SMs per thread block • If the number of blocks is smaller than that of SMs, the performance of the kernel becomes low. • Default block size • Cray compiler : 128 threads / block • Our compiler : 256 threads / block

  25. NPB CG 9.7x 2.1x 0.66x 0.74x the performance is lower than that of CPU and Cray compiler

  26. NPB CG • At class S, the performance of GPU is lower than that of CPU • Overheads are larger compared with kernel execution time • launching kernel functions • synchronization with device • data allocation / release / transfer • The overhead is larger than that of Cray compiler • large overhead of reduction • The performance of GPU kernels are better than that of Cray compiler

  27. Conclusion • We implemented a source-to-source OpenACC compiler for CUDA • C with OpenACCdirectives → C with CUDA API • Using Omni compiler infrastructure • In most case, the performance of GPU code by our compiler is higher than that of CPU single core • Speedup of up to 31 times at N-body • Our compiler makes use of CUDA backend successfully by source-to-source approach • the performance is often better than that of Cray compiler • There is room for performance improvement • using suitable grid size and block size • reducing overhead of synchronization and reduction

  28. Future Work • Optimization • tuning block size at compile time • reducing overhead from synchronization and reduction • Support the full set of directives for conforming to OpenACC specification in our compiler • We will release our compiler at next SC

More Related