openmp/libomptarget/deviceRTLs/nvptx/test/api/thread_limit.c

   1 // RUN: %compile-run-and-check
   2
   3 #include <omp.h>
   4 #include <stdio.h>
   5
   6 int main(int argc, char *argv[]) {
   7   int ThreadLimitL0 = -1, ThreadLimitL1 = -1, ThreadLimitL2 = -1;
   8
   9 #pragma omp declare reduction(unique64:int                                     \
  10                               : omp_out = (omp_in == 64 ? omp_in : omp_out))   \
  11     initializer(omp_priv = -1)
  12 #pragma omp declare reduction(unique32:int                                     \
  13                               : omp_out = (omp_in == 32 ? omp_in : omp_out))   \
  14     initializer(omp_priv = -1)
  15
  16   // Non-SPMD mode.
  17 #pragma omp target teams map(ThreadLimitL0, ThreadLimitL1, ThreadLimitL2)      \
  18     thread_limit(64) num_teams(1)
  19   {
  20     ThreadLimitL0 = omp_get_thread_limit();
  21 #pragma omp parallel reduction(unique64                                        \
  22                                : ThreadLimitL1, ThreadLimitL2) num_threads(32)
  23     {
  24       ThreadLimitL1 = omp_get_thread_limit();
  25 #pragma omp parallel reduction(unique64 : ThreadLimitL2)
  26       { ThreadLimitL2 = omp_get_thread_limit(); }
  27     }
  28   }
  29
  30   // CHECK: Non-SPMD ThreadLimitL0 = 64
  31   printf("Non-SPMD ThreadLimitL0 = %d\n", ThreadLimitL0);
  32   // CHECK: Non-SPMD ThreadLimitL1 = 64
  33   printf("Non-SPMD ThreadLimitL1 = %d\n", ThreadLimitL1);
  34   // CHECK: Non-SPMD ThreadLimitL2 = 64
  35   printf("Non-SPMD ThreadLimitL2 = %d\n", ThreadLimitL2);
  36
  37   // SPMD mode with full runtime
  38   ThreadLimitL1 = -1;
  39   ThreadLimitL2 = -1;
  40 #pragma omp target parallel reduction(unique32                                 \
  41                                       : ThreadLimitL1, ThreadLimitL2)          \
  42     num_threads(32)
  43   {
  44     ThreadLimitL1 = omp_get_thread_limit();
  45 #pragma omp parallel reduction(unique32 : ThreadLimitL2)
  46     { ThreadLimitL2 = omp_get_thread_limit(); }
  47   }
  48
  49   // CHECK: SPMD with full runtime ThreadLimitL1 = 32
  50   printf("SPMD with full runtime ThreadLimitL1 = %d\n", ThreadLimitL1);
  51   // CHECK: SPMD with full runtime ThreadLimitL2 = 32
  52   printf("SPMD with full runtime ThreadLimitL2 = %d\n", ThreadLimitL2);
  53
  54   // SPMD mode without runtime
  55   ThreadLimitL1 = -1;
  56   ThreadLimitL2 = -1;
  57 #pragma omp target parallel for reduction(unique32                             \
  58                                           : ThreadLimitL1, ThreadLimitL2)      \
  59     num_threads(32)
  60   for (int I = 0; I < 2; ++I) {
  61     ThreadLimitL1 = omp_get_thread_limit();
  62 #pragma omp parallel reduction(unique32 : ThreadLimitL2)
  63     { ThreadLimitL2 = omp_get_thread_limit(); }
  64   }
  65
  66   // CHECK: SPMD without runtime ThreadLimitL1 = 32
  67   printf("SPMD without runtime ThreadLimitL1 = %d\n", ThreadLimitL1);
  68   // CHECK: SPMD without runtime ThreadLimitL2 = 32
  69   printf("SPMD without runtime ThreadLimitL2 = %d\n", ThreadLimitL2);
  70
  71   return 0;
  72 }