clang/test/OpenMP/metadirective_device_arch_codegen.cpp

   1 // REQUIRES: amdgpu-registered-target
   2
   3 // RUN: %clang_cc1 -fopenmp -x c++ -w -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc
   4 // RUN: %clang_cc1 -fopenmp -x c++ -w -std=c++11 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -target-cpu gfx906 -o - | FileCheck %s
   5 // expected-no-diagnostics
   6
   7
   8 /*===-----------------------------------------------------------------------===
   9
  10 Inspired from SOLLVE tests:
  11  - 5.0/metadirective/test_metadirective_arch_is_nvidia.c
  12
  13
  14 ===------------------------------------------------------------------------===*/
  15
  16
  17 #define N 1024
  18
  19 int metadirective1() {
  20
  21    int v1[N], v2[N], v3[N];
  22
  23    int target_device_num, host_device_num, default_device;
  24    int errors = 0;
  25
  26    #pragma omp target map(to:v1,v2) map(from:v3, target_device_num) device(default_device)
  27    {
  28       #pragma omp metadirective \
  29                    when(device={arch("amdgcn")}: teams distribute parallel for) \
  30                    default(parallel for)
  31
  32          for (int i = 0; i < N; i++) {
  33             #pragma omp atomic write
  34             v3[i] = v1[i] * v2[i];
  35          }
  36    }
  37
  38    return errors;
  39 }
  40
  41 // CHECK: define weak_odr protected amdgpu_kernel void @[[METADIRECTIVE:.+metadirective1[a-z0-9_]+]]
  42 // CHECK: entry:
  43 // CHECK: %{{[0-9]}} = call i32 @__kmpc_target_init
  44 // CHECK: user_code.entry:
  45 // CHECK: call void @[[METADIRECTIVE]]_omp_outlined
  46 // CHECK-NOT: call void @__kmpc_parallel_51
  47 // CHECK: ret void
  48
  49
  50 // CHECK: define internal void @[[METADIRECTIVE]]_omp_outlined
  51 // CHECK: entry:
  52 // CHECK: call void @__kmpc_distribute_static_init
  53 // CHECK: omp.loop.exit:
  54 // CHECK: call void @__kmpc_distribute_static_fini
  55
  56
  57 // CHECK: define internal void @[[METADIRECTIVE]]_omp_outlined_omp_outlined
  58 // CHECK: entry:
  59 // CHECK: call void @__kmpc_for_static_init_4
  60 // CHECK: omp.inner.for.body:
  61 // CHECK: store atomic {{.*}} monotonic
  62 // CHECK: omp.loop.exit:
  63 // CHECK-NEXT: call void @__kmpc_for_static_fini
  64 // CHECK-NEXT: ret void
  65