| // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --prefix-filecheck-ir-name _ |
| // RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -fopenmp-version=52 -fopenmp-targets=nvptx64 -offload-device-only -x c -emit-llvm %s -o - | FileCheck %s |
| // expected-no-diagnostics |
| |
| void foo(); |
| |
| void masked() { |
| #pragma target |
| #pragma omp parallel masked |
| { |
| foo(); |
| } |
| } |
| |
| void maskedFilter() { |
| const int tid = 1; |
| #pragma target |
| #pragma omp parallel masked filter(tid) |
| { |
| foo(); |
| } |
| } |
| |
| void master() { |
| #pragma target |
| #pragma omp parallel master |
| { |
| foo(); |
| } |
| } |
| // CHECK-LABEL: define {{[^@]+}}@masked |
| // CHECK-SAME: () #[[ATTR0:[0-9]+]] { |
| // CHECK-NEXT: entry: |
| // CHECK-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1:[0-9]+]], i32 0, ptr @masked.omp_outlined) |
| // CHECK-NEXT: ret void |
| // |
| // |
| // CHECK-LABEL: define {{[^@]+}}@masked.omp_outlined |
| // CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] { |
| // CHECK-NEXT: entry: |
| // CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 |
| // CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 |
| // CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 |
| // CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 |
| // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8 |
| // CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 |
| // CHECK-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_masked(ptr @[[GLOB1]], i32 [[TMP1]], i32 0) |
| // CHECK-NEXT: [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0 |
| // CHECK-NEXT: br i1 [[TMP3]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]] |
| // CHECK: omp_if.then: |
| // CHECK-NEXT: call void (...) @foo() |
| // CHECK-NEXT: call void @__kmpc_end_masked(ptr @[[GLOB1]], i32 [[TMP1]]) |
| // CHECK-NEXT: br label [[OMP_IF_END]] |
| // CHECK: omp_if.end: |
| // CHECK-NEXT: ret void |
| // |
| // |
| // CHECK-LABEL: define {{[^@]+}}@maskedFilter |
| // CHECK-SAME: () #[[ATTR0]] { |
| // CHECK-NEXT: entry: |
| // CHECK-NEXT: [[TID:%.*]] = alloca i32, align 4 |
| // CHECK-NEXT: store i32 1, ptr [[TID]], align 4 |
| // CHECK-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @maskedFilter.omp_outlined) |
| // CHECK-NEXT: ret void |
| // |
| // |
| // CHECK-LABEL: define {{[^@]+}}@maskedFilter.omp_outlined |
| // CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] { |
| // CHECK-NEXT: entry: |
| // CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 |
| // CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 |
| // CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 |
| // CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 |
| // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8 |
| // CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 |
| // CHECK-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_masked(ptr @[[GLOB1]], i32 [[TMP1]], i32 1) |
| // CHECK-NEXT: [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0 |
| // CHECK-NEXT: br i1 [[TMP3]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]] |
| // CHECK: omp_if.then: |
| // CHECK-NEXT: call void (...) @foo() |
| // CHECK-NEXT: call void @__kmpc_end_masked(ptr @[[GLOB1]], i32 [[TMP1]]) |
| // CHECK-NEXT: br label [[OMP_IF_END]] |
| // CHECK: omp_if.end: |
| // CHECK-NEXT: ret void |
| // |
| // |
| // CHECK-LABEL: define {{[^@]+}}@master |
| // CHECK-SAME: () #[[ATTR0]] { |
| // CHECK-NEXT: entry: |
| // CHECK-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @master.omp_outlined) |
| // CHECK-NEXT: ret void |
| // |
| // |
| // CHECK-LABEL: define {{[^@]+}}@master.omp_outlined |
| // CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] { |
| // CHECK-NEXT: entry: |
| // CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 |
| // CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 |
| // CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 |
| // CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 |
| // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8 |
| // CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 |
| // CHECK-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_master(ptr @[[GLOB1]], i32 [[TMP1]]) |
| // CHECK-NEXT: [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0 |
| // CHECK-NEXT: br i1 [[TMP3]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]] |
| // CHECK: omp_if.then: |
| // CHECK-NEXT: call void (...) @foo() |
| // CHECK-NEXT: call void @__kmpc_end_master(ptr @[[GLOB1]], i32 [[TMP1]]) |
| // CHECK-NEXT: br label [[OMP_IF_END]] |
| // CHECK: omp_if.end: |
| // CHECK-NEXT: ret void |
| // |