162 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 10int *ptr;11void f1() {12 // &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM13 #pragma omp target map(ptr)14 ptr[1] = 6;15}16 17void f2() {18 // &ptr, &ptr[0], sizeof(ptr[0]), TO | FROM | PARAM | PTR_AND_OBJ19 #pragma omp target map(*ptr)20 ptr[1] = 6;21}22 23void f3() {24 // &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM25 // &ptr, &ptr[0], sizeof(ptr[0]), TO | FROM | PTR_AND_OBJ26 #pragma omp target map(ptr, *ptr)27 ptr[1] = 6;28}29 30void f4() {31 // &ptr, &ptr[0], sizeof(ptr[0]), TO | FROM | PTR_AND_OBJ | PARAM32 // &ptr, &ptr, sizeof(ptr), TO | FROM | PARAM33 #pragma omp target map(*ptr, ptr)34 ptr[2] = 8;35}36 37#endif38//.39// CHECK: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 8]40// CHECK: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x23]]]41// CHECK: @.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 4]42// CHECK: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]]43// CHECK: @.offload_sizes.3 = private unnamed_addr constant [2 x i64] [i64 8, i64 4]44// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [2 x i64] [i64 [[#0x23]], i64 [[#0x13]]]45// CHECK: @.offload_sizes.5 = private unnamed_addr constant [2 x i64] [i64 4, i64 8]46// CHECK: @.offload_maptypes.6 = private unnamed_addr constant [2 x i64] [i64 [[#0x33]], i64 [[#0x3]]]47//.48// CHECK-LABEL: define {{[^@]+}}@_Z2f1v49// CHECK-SAME: () #[[ATTR0:[0-9]+]] {50// CHECK: entry:51// CHECK: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 052// CHECK: store ptr @ptr, ptr [[TMP0]], align 853// CHECK: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 054// CHECK: store ptr @ptr, ptr [[TMP1]], align 855// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 056// CHECK: store ptr null, ptr [[TMP2]], align 857// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 058// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 059// CHECK: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 060//61//62// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f1v_l1363// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[PTR:%.*]]) #[[ATTR1:[0-9]+]] {64// CHECK: entry:65// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 866// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8, !nonnull [[META11:![0-9]+]], !align [[META12:![0-9]+]]67// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 868// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 169// CHECK: store i32 6, ptr [[ARRAYIDX]], align 470// CHECK: ret void71//72//73// CHECK-LABEL: define {{[^@]+}}@_Z2f2v74// CHECK-SAME: () #[[ATTR0]] {75// CHECK: entry:76// CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 877// CHECK: [[TMP1:%.*]] = load ptr, ptr @ptr, align 878// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 079// CHECK: store ptr @ptr, ptr [[TMP2]], align 880// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 081// CHECK: store ptr [[TMP1]], ptr [[TMP3]], align 882// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 083// CHECK: store ptr null, ptr [[TMP4]], align 884// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 085// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 086// CHECK: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 087//88//89// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f2v_l1990// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] {91// CHECK: entry:92// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 893// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 894// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 195// CHECK: store i32 6, ptr [[ARRAYIDX]], align 496// CHECK: ret void97//98//99// CHECK-LABEL: define {{[^@]+}}@_Z2f3v100// CHECK-SAME: () #[[ATTR0]] {101// CHECK: entry:102// CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8103// CHECK: [[TMP1:%.*]] = load ptr, ptr @ptr, align 8104// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0105// CHECK: store ptr @ptr, ptr [[TMP2]], align 8106// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0107// CHECK: store ptr @ptr, ptr [[TMP3]], align 8108// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0109// CHECK: store ptr null, ptr [[TMP4]], align 8110// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1111// CHECK: store ptr @ptr, ptr [[TMP5]], align 8112// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1113// CHECK: store ptr [[TMP1]], ptr [[TMP6]], align 8114// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1115// CHECK: store ptr null, ptr [[TMP7]], align 8116// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0117// CHECK: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0118// CHECK: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0119//120//121// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f3v_l26122// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] {123// CHECK: entry:124// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8125// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8126// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 1127// CHECK: store i32 6, ptr [[ARRAYIDX]], align 4128// CHECK: ret void129//130//131// CHECK-LABEL: define {{[^@]+}}@_Z2f4v132// CHECK-SAME: () #[[ATTR0]] {133// CHECK: entry:134// CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8135// CHECK: [[TMP1:%.*]] = load ptr, ptr @ptr, align 8136// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0137// CHECK: store ptr @ptr, ptr [[TMP2]], align 8138// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0139// CHECK: store ptr [[TMP1]], ptr [[TMP3]], align 8140// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0141// CHECK: store ptr null, ptr [[TMP4]], align 8142// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1143// CHECK: store ptr @ptr, ptr [[TMP5]], align 8144// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1145// CHECK: store ptr @ptr, ptr [[TMP6]], align 8146// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1147// CHECK: store ptr null, ptr [[TMP7]], align 8148// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0149// CHECK: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0150// CHECK: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0151//152//153// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f4v_l33154// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] {155// CHECK: entry:156// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8157// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8158// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2159// CHECK: store i32 8, ptr [[ARRAYIDX]], align 4160// CHECK: ret void161//162