openmp/libomptarget/deviceRTLs/nvptx/test/api/max_threads.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 MaxThreadsL1 = -1, MaxThreadsL2 = -1;
   8
   9 #pragma omp declare reduction(unique:int                                       \
  10                               : omp_out = (omp_in == 1 ? omp_in : omp_out))    \
  11     initializer(omp_priv = -1)
  12
  13   // Non-SPMD mode.
  14 #pragma omp target teams map(MaxThreadsL1, MaxThreadsL2) thread_limit(32)      \
  15     num_teams(1)
  16   {
  17     MaxThreadsL1 = omp_get_max_threads();
  18 #pragma omp parallel reduction(unique : MaxThreadsL2)
  19     { MaxThreadsL2 = omp_get_max_threads(); }
  20   }
  21
  22   //FIXME: This Non-SPMD kernel will have 32 active threads due to
  23   //       thread_limit. However, Non-SPMD MaxThreadsL1 is the total number of
  24   //       threads in block (64 in this case), which translates to worker
  25   //       threads + WARP_SIZE for Non-SPMD kernels and worker threads for SPMD
  26   //       kernels. According to the spec, omp_get_max_threads must return the
  27   //       max active threads possible between the two kernel types.
  28
  29   // CHECK: Non-SPMD MaxThreadsL1 = 64
  30   printf("Non-SPMD MaxThreadsL1 = %d\n", MaxThreadsL1);
  31   // CHECK: Non-SPMD MaxThreadsL2 = 1
  32   printf("Non-SPMD MaxThreadsL2 = %d\n", MaxThreadsL2);
  33
  34   // SPMD mode with full runtime
  35   MaxThreadsL2 = -1;
  36 #pragma omp target parallel reduction(unique : MaxThreadsL2)
  37   { MaxThreadsL2 = omp_get_max_threads(); }
  38
  39   // CHECK: SPMD with full runtime MaxThreadsL2 = 1
  40   printf("SPMD with full runtime MaxThreadsL2 = %d\n", MaxThreadsL2);
  41
  42   // SPMD mode without runtime
  43   MaxThreadsL2 = -1;
  44 #pragma omp target parallel for reduction(unique : MaxThreadsL2)
  45   for (int I = 0; I < 2; ++I) {
  46     MaxThreadsL2 = omp_get_max_threads();
  47   }
  48
  49   // CHECK: SPMD without runtime MaxThreadsL2 = 1
  50   printf("SPMD without runtime MaxThreadsL2 = %d\n", MaxThreadsL2);
  51
  52   return 0;
  53 }