brintos

brintos / llvm-project-archived public Read only

0
0
Text · 1.4 KiB · 8b7b999 Raw
49 lines · plain
1Rewriting generic-mode kernel with a customized state machine. [OMP131]2=======================================================================3 4.. _omp131:5 6This optimization remark indicates that a generic-mode kernel on the device was7specialized for the given target region. When offloading in generic-mode, a8state machine is required to schedule the work between the parallel worker9threads. This optimization specializes the state machine in cases where there is10a known number of parallel regions inside the kernel. A much simpler state11machine can be used if it is known that there is no nested parallelism and the12number of regions to schedule is a static amount.13 14Examples15--------16 17This optimization should occur on any generic-mode kernel that has visibility on18all parallel regions, but cannot be moved to SPMD-mode and has no nested19parallelism.20 21.. code-block:: c++22 23   #pragma omp declare target24   int TID;25   #pragma omp end declare target26 27   void foo() {28   #pragma omp target29   {30    TID = omp_get_thread_num();31    #pragma omp parallel32    {33      work();34    }35   }36   }37 38.. code-block:: console39 40   $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass=openmp-opt omp131.cpp41   omp131.cpp:8:1: remark: Rewriting generic-mode kernel with a customized state machine. [OMP131]42   #pragma omp target43   ^44 45Diagnostic Scope46----------------47 48OpenMP target offloading optimization remark.49