1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" 2 // RUN: %clang_cc1 -no-opaque-pointers -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -fopenmp-offload-mandatory -emit-llvm %s -o - | FileCheck %s --check-prefix=MANDATORY 3 // expected-no-diagnostics 4 5 void foo() {} 6 #pragma omp declare target(foo) 7 8 void bar() {} 9 #pragma omp declare target device_type(nohost) to(bar) 10 11 void host() { 12 #pragma omp target 13 { bar(); } 14 } 15 16 void host_if(bool cond) { 17 #pragma omp target if(cond) 18 { bar(); } 19 } 20 21 void host_dev(int device) { 22 #pragma omp target device(device) 23 { bar(); } 24 } 25 // MANDATORY-LABEL: define {{[^@]+}}@_Z3foov 26 // MANDATORY-SAME: () #[[ATTR0:[0-9]+]] { 27 // MANDATORY-NEXT: entry: 28 // MANDATORY-NEXT: ret void 29 // 30 // 31 // MANDATORY-LABEL: define {{[^@]+}}@_Z4hostv 32 // MANDATORY-SAME: () #[[ATTR0]] { 33 // MANDATORY-NEXT: entry: 34 // MANDATORY-NEXT: [[TMP0:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4hostv_l12.region_id, i32 0, i8** null, i8** null, i64* null, i64* null, i8** null, i8** null) 35 // MANDATORY-NEXT: [[TMP1:%.*]] = icmp ne i32 [[TMP0]], 0 36 // MANDATORY-NEXT: br i1 [[TMP1]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 37 // MANDATORY: omp_offload.failed: 38 // MANDATORY-NEXT: unreachable 39 // MANDATORY: omp_offload.cont: 40 // MANDATORY-NEXT: ret void 41 // 42 // 43 // MANDATORY-LABEL: define {{[^@]+}}@_Z7host_ifb 44 // MANDATORY-SAME: (i1 noundef zeroext [[COND:%.*]]) #[[ATTR0]] { 45 // MANDATORY-NEXT: entry: 46 // MANDATORY-NEXT: [[COND_ADDR:%.*]] = alloca i8, align 1 47 // MANDATORY-NEXT: [[FROMBOOL:%.*]] = zext i1 [[COND]] to i8 48 // MANDATORY-NEXT: store i8 [[FROMBOOL]], i8* [[COND_ADDR]], align 1 49 // MANDATORY-NEXT: [[TMP0:%.*]] = load i8, i8* [[COND_ADDR]], align 1 50 // MANDATORY-NEXT: [[TOBOOL:%.*]] = trunc i8 [[TMP0]] to i1 51 // MANDATORY-NEXT: br i1 [[TOBOOL]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]] 52 // MANDATORY: omp_if.then: 53 // MANDATORY-NEXT: [[TMP1:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z7host_ifb_l17.region_id, i32 0, i8** null, i8** null, i64* null, i64* null, i8** null, i8** null) 54 // MANDATORY-NEXT: [[TMP2:%.*]] = icmp ne i32 [[TMP1]], 0 55 // MANDATORY-NEXT: br i1 [[TMP2]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 56 // MANDATORY: omp_offload.failed: 57 // MANDATORY-NEXT: unreachable 58 // MANDATORY: omp_offload.cont: 59 // MANDATORY-NEXT: br label [[OMP_IF_END:%.*]] 60 // MANDATORY: omp_if.else: 61 // MANDATORY-NEXT: unreachable 62 // MANDATORY: omp_if.end: 63 // MANDATORY-NEXT: ret void 64 // 65 // 66 // MANDATORY-LABEL: define {{[^@]+}}@_Z8host_devi 67 // MANDATORY-SAME: (i32 noundef signext [[DEVICE:%.*]]) #[[ATTR0]] { 68 // MANDATORY-NEXT: entry: 69 // MANDATORY-NEXT: [[DEVICE_ADDR:%.*]] = alloca i32, align 4 70 // MANDATORY-NEXT: [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4 71 // MANDATORY-NEXT: store i32 [[DEVICE]], i32* [[DEVICE_ADDR]], align 4 72 // MANDATORY-NEXT: [[TMP0:%.*]] = load i32, i32* [[DEVICE_ADDR]], align 4 73 // MANDATORY-NEXT: store i32 [[TMP0]], i32* [[DOTCAPTURE_EXPR_]], align 4 74 // MANDATORY-NEXT: [[TMP1:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_]], align 4 75 // MANDATORY-NEXT: [[TMP2:%.*]] = sext i32 [[TMP1]] to i64 76 // MANDATORY-NEXT: [[TMP3:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 [[TMP2]], i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z8host_devi_l22.region_id, i32 0, i8** null, i8** null, i64* null, i64* null, i8** null, i8** null) 77 // MANDATORY-NEXT: [[TMP4:%.*]] = icmp ne i32 [[TMP3]], 0 78 // MANDATORY-NEXT: br i1 [[TMP4]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 79 // MANDATORY: omp_offload.failed: 80 // MANDATORY-NEXT: unreachable 81 // MANDATORY: omp_offload.cont: 82 // MANDATORY-NEXT: ret void 83 // 84 // 85 // MANDATORY-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 86 // MANDATORY-SAME: () #[[ATTR3:[0-9]+]] { 87 // MANDATORY-NEXT: entry: 88 // MANDATORY-NEXT: call void @__tgt_register_requires(i64 1) 89 // MANDATORY-NEXT: ret void 90 // 91