168 lines · cpp
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --check-globals --filter-out-after "getelem.*kernel" --filter-out "= alloca.*" --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _ --global-value-regex "\.offload_.*" --global-hex-value-regex ".offload_maptypes.*"2// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s3// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s4// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s5 6// expected-no-diagnostics7#ifndef HEADER8#define HEADER9 10void f1() {11 int *ptr;12 // &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM13 #pragma omp target map(ptr)14 ptr[1] = 6;15}16 17void f2() {18 int *ptr;19 // &ptr[0], &ptr[0], sizeof(ptr[0]), TO | FROM | PARAM20 #pragma omp target map(*ptr)21 ptr[1] = 6;22}23 24void f3() {25 int *ptr;26 // &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM27 // &ptr[0], &ptr[0], sizeof(ptr[0]), TO | FROM28 #pragma omp target map(ptr, *ptr)29 ptr[1] = 6;30}31 32void f4() {33 int *ptr;34 // &ptr[0], &ptr[0], sizeof(ptr[0]), TO | FROM35 // &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM36 #pragma omp target map(*ptr, ptr)37 ptr[2] = 8;38}39 40#endif41//.42// CHECK: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 8]43// CHECK: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x23]]]44// CHECK: @.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 4]45// CHECK: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 [[#0x23]]]46// CHECK: @.offload_sizes.3 = private unnamed_addr constant [2 x i64] [i64 8, i64 4]47// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [2 x i64] [i64 [[#0x23]], i64 [[#0x3]]]48// CHECK: @.offload_sizes.5 = private unnamed_addr constant [2 x i64] [i64 4, i64 8]49// CHECK: @.offload_maptypes.6 = private unnamed_addr constant [2 x i64] [i64 [[#0x23]], i64 [[#0x3]]]50//.51// CHECK-LABEL: define {{[^@]+}}@_Z2f1v52// CHECK-SAME: () #[[ATTR0:[0-9]+]] {53// CHECK: entry:54// CHECK: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 055// CHECK: store ptr [[PTR:%.*]], ptr [[TMP0]], align 856// CHECK: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 057// CHECK: store ptr [[PTR]], ptr [[TMP1]], align 858// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 059// CHECK: store ptr null, ptr [[TMP2]], align 860// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 061// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 062// CHECK: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 063//64//65// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f1v_l1366// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[PTR:%.*]]) #[[ATTR1:[0-9]+]] {67// CHECK: entry:68// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 869// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8, !nonnull [[META11:![0-9]+]], !align [[META12:![0-9]+]]70// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 871// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 172// CHECK: store i32 6, ptr [[ARRAYIDX]], align 473// CHECK: ret void74//75//76// CHECK-LABEL: define {{[^@]+}}@_Z2f2v77// CHECK-SAME: () #[[ATTR0]] {78// CHECK: entry:79// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR:%.*]], align 880// CHECK: [[TMP1:%.*]] = load ptr, ptr [[PTR]], align 881// CHECK: [[TMP2:%.*]] = load ptr, ptr [[PTR]], align 882// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 083// CHECK: store ptr [[TMP1]], ptr [[TMP3]], align 884// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 085// CHECK: store ptr [[TMP2]], ptr [[TMP4]], align 886// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 087// CHECK: store ptr null, ptr [[TMP5]], align 888// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 089// CHECK: [[TMP7:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 090// CHECK: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 091//92//93// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f2v_l2094// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] {95// CHECK: entry:96// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 897// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 898// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 199// CHECK: store i32 6, ptr [[ARRAYIDX]], align 4100// CHECK: ret void101//102//103// CHECK-LABEL: define {{[^@]+}}@_Z2f3v104// CHECK-SAME: () #[[ATTR0]] {105// CHECK: entry:106// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR:%.*]], align 8107// CHECK: [[TMP1:%.*]] = load ptr, ptr [[PTR]], align 8108// CHECK: [[TMP2:%.*]] = load ptr, ptr [[PTR]], align 8109// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0110// CHECK: store ptr [[PTR]], ptr [[TMP3]], align 8111// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0112// CHECK: store ptr [[PTR]], ptr [[TMP4]], align 8113// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0114// CHECK: store ptr null, ptr [[TMP5]], align 8115// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1116// CHECK: store ptr [[TMP1]], ptr [[TMP6]], align 8117// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1118// CHECK: store ptr [[TMP2]], ptr [[TMP7]], align 8119// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1120// CHECK: store ptr null, ptr [[TMP8]], align 8121// CHECK: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0122// CHECK: [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0123// CHECK: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0124//125//126// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f3v_l28127// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] {128// CHECK: entry:129// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8130// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8131// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 1132// CHECK: store i32 6, ptr [[ARRAYIDX]], align 4133// CHECK: ret void134//135//136// CHECK-LABEL: define {{[^@]+}}@_Z2f4v137// CHECK-SAME: () #[[ATTR0]] {138// CHECK: entry:139// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR:%.*]], align 8140// CHECK: [[TMP1:%.*]] = load ptr, ptr [[PTR]], align 8141// CHECK: [[TMP2:%.*]] = load ptr, ptr [[PTR]], align 8142// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0143// CHECK: store ptr [[TMP1]], ptr [[TMP3]], align 8144// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0145// CHECK: store ptr [[TMP2]], ptr [[TMP4]], align 8146// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0147// CHECK: store ptr null, ptr [[TMP5]], align 8148// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1149// CHECK: store ptr [[PTR]], ptr [[TMP6]], align 8150// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1151// CHECK: store ptr [[PTR]], ptr [[TMP7]], align 8152// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1153// CHECK: store ptr null, ptr [[TMP8]], align 8154// CHECK: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0155// CHECK: [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0156// CHECK: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0157//158//159// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f4v_l36160// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] {161// CHECK: entry:162// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8163// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8164// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2165// CHECK: store i32 8, ptr [[ARRAYIDX]], align 4166// CHECK: ret void167//168