test/OpenMP/parallel_masked_target.cpp

// 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
//