90 lines · plain
1.. _omp112:2 3Found thread data sharing on the GPU. Expect degraded performance due to data globalization. [OMP112]4=====================================================================================================5 6This missed remark indicates that a globalized value was found on the target7device that was not either replaced with stack memory by :ref:`OMP110 <omp110>`8or shared memory by :ref:`OMP111 <omp111>`. Globalization that has not been9removed will need to be handled by the runtime and will significantly impact10performance.11 12The OpenMP standard requires that threads are able to share their data between13each-other. However, this is not true by default when offloading to a target14device such as a GPU. Threads on a GPU cannot shared their data unless it is15first placed in global or shared memory. In order to create standards complaint16code, the Clang compiler will globalize any variables that could potentially be17shared between the threads. In the majority of cases, globalized variables can18either be returns to a thread-local stack, or pushed to shared memory. However,19in a few cases it is necessary and will cause a performance penalty.20 21Examples22--------23 24This example shows legitimate data sharing on the device. It is a convoluted25example, but is completely complaint with the OpenMP standard. If globalization26was not added this would result in different results on different target27devices.28 29.. code-block:: c++30 31 #include <omp.h>32 #include <cstdio>33 34 #pragma omp declare target35 static int *p;36 #pragma omp end declare target37 38 void foo() {39 int x = omp_get_thread_num();40 if (omp_get_thread_num() == 1)41 p = &x;42 43 #pragma omp barrier44 45 printf ("Thread %d: %d\n", omp_get_thread_num(), *p);46 }47 48 int main() {49 #pragma omp target parallel50 foo();51 }52 53.. code-block:: console54 55 $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O1 -Rpass-missed=openmp-opt omp112.cpp56 omp112.cpp:9:7: remark: Found thread data sharing on the GPU. Expect degraded performance57 due to data globalization. [OMP112] [-Rpass-missed=openmp-opt]58 int x = omp_get_thread_num();59 ^60 61A less convoluted example globalization that cannot be removed occurs when62calling functions that aren't visible from the current translation unit.63 64.. code-block:: c++65 66 extern void use(int *x);67 68 void foo() {69 int x;70 use(&x);71 }72 73 int main() {74 #pragma omp target parallel75 foo();76 }77 78.. code-block:: console79 80 $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O1 -Rpass-missed=openmp-opt omp112.cpp81 omp112.cpp:4:7: remark: Found thread data sharing on the GPU. Expect degraded performance82 due to data globalization. [OMP112] [-Rpass-missed=openmp-opt]83 int x;84 ^85 86Diagnostic Scope87----------------88 89OpenMP target offloading missed remark.90