clang/test/OpenMP/thread_limit_nvptx.c

   1 // Test target codegen - host bc file has to be created first.
   2 //
   3 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
   4 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s
   5 // expected-no-diagnostics
   6 #ifndef HEADER
   7 #define HEADER
   8
   9 void foo(int N) {
  10 // CHECK: l11, !"maxntidx", i32 128}
  11 #pragma omp target teams distribute parallel for simd
  12   for (int i = 0; i < N; ++i)
  13     ;
  14 // CHECK: l15, !"maxntidx", i32 4}
  15 #pragma omp target teams distribute parallel for simd thread_limit(4)
  16   for (int i = 0; i < N; ++i)
  17     ;
  18 // CHECK-NOT: l21, !"maxntidx", i32 128}
  19 // CHECK: l21, !"maxntidx", i32 42}
  20 // CHECK-NOT: l21, !"maxntidx", i32 128}
  21 #pragma omp target teams distribute parallel for simd ompx_attribute(__attribute__((launch_bounds(42, 42))))
  22   for (int i = 0; i < N; ++i)
  23     ;
  24 // CHECK-NOT: l27, !"maxntidx", i32 42}
  25 // CHECK: l27, !"maxntidx", i32 22}
  26 // CHECK-NOT: l27, !"maxntidx", i32 42}
  27 #pragma omp target teams distribute parallel for simd ompx_attribute(__attribute__((launch_bounds(42, 42)))) num_threads(22)
  28   for (int i = 0; i < N; ++i)
  29     ;
  30 }
  31
  32 #endif
  33