[Openmp-dev] OpenMP GPU Target Offload in Clang

[Qiongsi Wu via Openmp-dev]

Hi Joachim!


Thanks for your help! I missed the cmake flags for OpenMP targets when building. Additionally, I found that my libelf was not installed properly. After rebuilding clang and removing the reduction code, the code offload was successful.


The debugging code you suggested works like a charm!


Sincerely,

Qiongsi

________________________________
From: Joachim Protze <protze.joachim at gmail.com>
Sent: Tuesday, August 21, 2018 12:37:02 PM
To: Qiongsi Wu; OpenMP-dev
Subject: Re: [Openmp-dev] OpenMP GPU Target Offload in Clang

Hi Qiongsi,

On 08/21/2018 05:28 PM, Qiongsi Wu via Openmp-dev wrote:
> Hi, OpenMP dev community!
>
>
> Recently I tried setting up the OpenMP benchmarks for SPEC ACCEL and
> test it with clang, but I ran into several difficulties.
>
>
> The core of the issue is that I was not able to get the workload onto
> the GPUs.  I wrote the following small test
>
>
> //////////////////////////////////////////////////////////////////////////////////////////////////////////////
>
> #define DATATYPE unsigned long long
>
> /*gpu offload openmp*/
> DATATYPE reduce_gpu_omp(DATATYPE *arr, size_t size) {
>      DATATYPE result = IDENTITY;
> #pragma omp target data map(tofrom:arr[:size]) map(tofrom:result)
>      {
> #pragma omp target teams distribute parallel for reduction(+:result)
> schedule(static, 1)
>          for (size_t i = 0; i < size; i++) {
>              result += arr[i];
>          }
>      }
>      return result;
> }
>
> //////////////////////////////////////////////////////////////////////////////////////////////////////////////
>

When compiling your code, but leaving out the reduction, I can execute
this on a GPU. With the reduction the code seems to hang for me.
(Posting a full compile-able example next time would be preferred!)

This is how I compiled:

clang -O3 -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -o reduce reduce.c

To see whether the code is actually executed on the device, you can add
this to the loop for debugging:

if (i==0) printf("omp_is_initial_device=%i\n", omp_is_initial_device());

> And compiled that with clang trunk with the following commands:
>
> clang -O3 -fopenmp -omptargets=nvptx64sm_35-nvidia-linux -Wall -o reduce
> reduce.c
> clang -O3 -fopenmp -omptargets=nvptx64sm_35-nvidia-linux-cuda -Wall -o
> reduce reduce.c
> clang -O3 -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -Wall -o reduce
> reduce.c
>
> The offloading to GPU was unsuccessful for all these commands. That
> said, the CPU load did go up when the kernel above was run,  so the
> offloading did happen, but the computation was offloaded to the CPU, not
> the GPU.
>
> My speculation is that I missed some steps setting up the
> compiler/libraries and the offloading did not happen correctly. Or it
> could be the fact that reductions were not supported across teams (as
> stated here https://clang.llvm.org/docs/OpenMPSupport.html).
>
> In the end, I would like to ask two questions:
>
>  1. What is a good candidate of llvm based compiler to test OpenMP GPU
>     offloading?  Should clang-ykt be used instead of clang trunk?

I used clang trunk, compiled like:

cmake -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=$INSTALL  \
       -DCLANG_OPENMP_NVPTX_DEFAULT_ARCH=sm_60 \
       -DLIBOMPTARGET_NVPTX_COMPUTE_CAPABILITIES=35,60 \
       $SRC


>  2. What is the recommended procedure for compiler and linker flags to
>     build programs with GPU offloading? Maybe I am not searching
>     correctly, but I was not able to find a documentation on how that is
>     supposed to be done. Additionally, will the compiler show some
>     warning if offloading to GPU is unsuccessful?

Successful / unsuccessful is a runtime decision. You will get an error
and execution aborts, once this patch is submitted:

https://reviews.llvm.org/D50522

Best
Joachim

>
>
> Thanks for your help!
>
> Sincerely,
> Qiongsi
>
>
>
> _______________________________________________
> Openmp-dev mailing list
> Openmp-dev at lists.llvm.org
> http://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev
>

-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/openmp-dev/attachments/20180821/96b8597e/attachment.html>