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 #ifndef HEADER 8 #define HEADER 9 amdgcn_device_isa_selected()10int amdgcn_device_isa_selected() { 11 int threadCount = 0; 12 13 #pragma omp target map(tofrom \ 14 : threadCount) 15 { 16 #pragma omp metadirective \ 17 when(device = {isa("dpp")} \ 18 : parallel) default(single) 19 threadCount++; 20 } 21 22 return threadCount; 23 } 24 25 // CHECK: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}amdgcn_device_isa_selected 26 // CHECK: user_code.entry: 27 // CHECK: call void @__kmpc_parallel_51 28 // CHECK-NOT: call i32 @__kmpc_single 29 // CHECK: ret void 30 amdgcn_device_isa_not_selected()31int amdgcn_device_isa_not_selected() { 32 int threadCount = 0; 33 34 #pragma omp target map(tofrom \ 35 : threadCount) 36 { 37 #pragma omp metadirective \ 38 when(device = {isa("sse")} \ 39 : parallel) \ 40 when(device = {isa("another-unsupported-gpu-feature")} \ 41 : parallel) default(single) 42 threadCount++; 43 } 44 45 return threadCount; 46 } 47 // CHECK: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}amdgcn_device_isa_not_selected 48 // CHECK: user_code.entry: 49 // CHECK: call i32 @__kmpc_single 50 // CHECK-NOT: call void @__kmpc_parallel_51 51 // CHECK: ret void 52 53 #endif 54