forked from OSchip/llvm-project
2233 lines
160 KiB
C++
2233 lines
160 KiB
C++
// 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]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _
|
|
// REQUIRES: powerpc-registered-target
|
|
// REQUIRES: nvptx-registered-target
|
|
|
|
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK1
|
|
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
|
|
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK2
|
|
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t
|
|
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - | FileCheck %s --check-prefix=CHECK3
|
|
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - | FileCheck %s --check-prefix=CHECK4
|
|
|
|
// expected-no-diagnostics
|
|
#ifndef HEADER
|
|
#define HEADER
|
|
|
|
template <typename T>
|
|
int foo(const T &t) {
|
|
#pragma omp target parallel
|
|
t();
|
|
return 0;
|
|
}
|
|
|
|
struct S {
|
|
int a = 15;
|
|
int foo() {
|
|
auto &&L = [&]() { return a; };
|
|
#pragma omp target
|
|
L();
|
|
#pragma omp target parallel
|
|
L();
|
|
return a + ::foo(L);
|
|
}
|
|
} s;
|
|
|
|
int main(int argc, char **argv) {
|
|
int &b = argc;
|
|
int &&c = 1;
|
|
int *d = &argc;
|
|
int a;
|
|
auto &&L = [&]() { return argc + b + c + reinterpret_cast<long int>(d) + a; };
|
|
#pragma omp target firstprivate(argc) map(to : a)
|
|
L();
|
|
#pragma omp target parallel
|
|
L();
|
|
return argc + s.foo();
|
|
}
|
|
|
|
#endif // HEADER
|
|
// CHECK1-LABEL: define {{[^@]+}}@main
|
|
// CHECK1-SAME: (i32 signext [[ARGC:%.*]], i8** [[ARGV:%.*]]) #[[ATTR0:[0-9]+]] {
|
|
// CHECK1-NEXT: entry:
|
|
// CHECK1-NEXT: [[RETVAL:%.*]] = alloca i32, align 4
|
|
// CHECK1-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4
|
|
// CHECK1-NEXT: [[ARGV_ADDR:%.*]] = alloca i8**, align 8
|
|
// CHECK1-NEXT: [[B:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[C:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[REF_TMP:%.*]] = alloca i32, align 4
|
|
// CHECK1-NEXT: [[D:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[A:%.*]] = alloca i32, align 4
|
|
// CHECK1-NEXT: [[L:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK1-NEXT: [[REF_TMP1:%.*]] = alloca [[CLASS_ANON:%.*]], align 8
|
|
// CHECK1-NEXT: [[TMP:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[_TMP2:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[_TMP3:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK1-NEXT: [[ARGC_CASTED:%.*]] = alloca i64, align 8
|
|
// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [11 x i8*], align 8
|
|
// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [11 x i8*], align 8
|
|
// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [11 x i8*], align 8
|
|
// CHECK1-NEXT: [[_TMP4:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[_TMP5:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[_TMP6:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS7:%.*]] = alloca [11 x i8*], align 8
|
|
// CHECK1-NEXT: [[DOTOFFLOAD_PTRS8:%.*]] = alloca [11 x i8*], align 8
|
|
// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS9:%.*]] = alloca [11 x i8*], align 8
|
|
// CHECK1-NEXT: store i32 0, i32* [[RETVAL]], align 4
|
|
// CHECK1-NEXT: store i32 [[ARGC]], i32* [[ARGC_ADDR]], align 4
|
|
// CHECK1-NEXT: store i8** [[ARGV]], i8*** [[ARGV_ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[ARGC_ADDR]], i32** [[B]], align 8
|
|
// CHECK1-NEXT: store i32 1, i32* [[REF_TMP]], align 4
|
|
// CHECK1-NEXT: store i32* [[REF_TMP]], i32** [[C]], align 8
|
|
// CHECK1-NEXT: store i32* [[ARGC_ADDR]], i32** [[D]], align 8
|
|
// CHECK1-NEXT: [[TMP0:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[REF_TMP1]], i32 0, i32 0
|
|
// CHECK1-NEXT: store i32* [[ARGC_ADDR]], i32** [[TMP0]], align 8
|
|
// CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[REF_TMP1]], i32 0, i32 1
|
|
// CHECK1-NEXT: [[TMP2:%.*]] = load i32*, i32** [[B]], align 8
|
|
// CHECK1-NEXT: store i32* [[TMP2]], i32** [[TMP1]], align 8
|
|
// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[REF_TMP1]], i32 0, i32 2
|
|
// CHECK1-NEXT: [[TMP4:%.*]] = load i32*, i32** [[C]], align 8
|
|
// CHECK1-NEXT: store i32* [[TMP4]], i32** [[TMP3]], align 8
|
|
// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[REF_TMP1]], i32 0, i32 3
|
|
// CHECK1-NEXT: store i32** [[D]], i32*** [[TMP5]], align 8
|
|
// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[REF_TMP1]], i32 0, i32 4
|
|
// CHECK1-NEXT: store i32* [[A]], i32** [[TMP6]], align 8
|
|
// CHECK1-NEXT: store %class.anon* [[REF_TMP1]], %class.anon** [[L]], align 8
|
|
// CHECK1-NEXT: [[TMP7:%.*]] = load i32*, i32** [[B]], align 8
|
|
// CHECK1-NEXT: store i32* [[TMP7]], i32** [[TMP]], align 8
|
|
// CHECK1-NEXT: [[TMP8:%.*]] = load i32*, i32** [[C]], align 8
|
|
// CHECK1-NEXT: store i32* [[TMP8]], i32** [[_TMP2]], align 8
|
|
// CHECK1-NEXT: [[TMP9:%.*]] = load %class.anon*, %class.anon** [[L]], align 8
|
|
// CHECK1-NEXT: store %class.anon* [[TMP9]], %class.anon** [[_TMP3]], align 8
|
|
// CHECK1-NEXT: [[TMP10:%.*]] = load i32, i32* [[ARGC_ADDR]], align 4
|
|
// CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[ARGC_CASTED]] to i32*
|
|
// CHECK1-NEXT: store i32 [[TMP10]], i32* [[CONV]], align 4
|
|
// CHECK1-NEXT: [[TMP11:%.*]] = load i64, i64* [[ARGC_CASTED]], align 8
|
|
// CHECK1-NEXT: [[TMP12:%.*]] = load i32*, i32** [[TMP]], align 8
|
|
// CHECK1-NEXT: [[TMP13:%.*]] = load i32*, i32** [[_TMP2]], align 8
|
|
// CHECK1-NEXT: [[TMP14:%.*]] = load i32*, i32** [[D]], align 8
|
|
// CHECK1-NEXT: [[TMP15:%.*]] = load %class.anon*, %class.anon** [[_TMP3]], align 8
|
|
// CHECK1-NEXT: [[TMP16:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP15]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP15]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP18:%.*]] = load i32*, i32** [[TMP17]], align 8
|
|
// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP15]], i32 0, i32 1
|
|
// CHECK1-NEXT: [[TMP20:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP15]], i32 0, i32 1
|
|
// CHECK1-NEXT: [[TMP21:%.*]] = load i32*, i32** [[TMP20]], align 8
|
|
// CHECK1-NEXT: [[TMP22:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP15]], i32 0, i32 2
|
|
// CHECK1-NEXT: [[TMP23:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP15]], i32 0, i32 2
|
|
// CHECK1-NEXT: [[TMP24:%.*]] = load i32*, i32** [[TMP23]], align 8
|
|
// CHECK1-NEXT: [[TMP25:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP15]], i32 0, i32 3
|
|
// CHECK1-NEXT: [[TMP26:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP15]], i32 0, i32 3
|
|
// CHECK1-NEXT: [[TMP27:%.*]] = load i32**, i32*** [[TMP26]], align 8
|
|
// CHECK1-NEXT: [[TMP28:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP15]], i32 0, i32 4
|
|
// CHECK1-NEXT: [[TMP29:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP15]], i32 0, i32 4
|
|
// CHECK1-NEXT: [[TMP30:%.*]] = load i32*, i32** [[TMP29]], align 8
|
|
// CHECK1-NEXT: [[TMP31:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP32:%.*]] = bitcast i8** [[TMP31]] to i64*
|
|
// CHECK1-NEXT: store i64 [[TMP11]], i64* [[TMP32]], align 8
|
|
// CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP34:%.*]] = bitcast i8** [[TMP33]] to i64*
|
|
// CHECK1-NEXT: store i64 [[TMP11]], i64* [[TMP34]], align 8
|
|
// CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP35]], align 8
|
|
// CHECK1-NEXT: [[TMP36:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
|
|
// CHECK1-NEXT: [[TMP37:%.*]] = bitcast i8** [[TMP36]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP12]], i32** [[TMP37]], align 8
|
|
// CHECK1-NEXT: [[TMP38:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1
|
|
// CHECK1-NEXT: [[TMP39:%.*]] = bitcast i8** [[TMP38]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP12]], i32** [[TMP39]], align 8
|
|
// CHECK1-NEXT: [[TMP40:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP40]], align 8
|
|
// CHECK1-NEXT: [[TMP41:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2
|
|
// CHECK1-NEXT: [[TMP42:%.*]] = bitcast i8** [[TMP41]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP13]], i32** [[TMP42]], align 8
|
|
// CHECK1-NEXT: [[TMP43:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2
|
|
// CHECK1-NEXT: [[TMP44:%.*]] = bitcast i8** [[TMP43]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP13]], i32** [[TMP44]], align 8
|
|
// CHECK1-NEXT: [[TMP45:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP45]], align 8
|
|
// CHECK1-NEXT: [[TMP46:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3
|
|
// CHECK1-NEXT: [[TMP47:%.*]] = bitcast i8** [[TMP46]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP14]], i32** [[TMP47]], align 8
|
|
// CHECK1-NEXT: [[TMP48:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 3
|
|
// CHECK1-NEXT: [[TMP49:%.*]] = bitcast i8** [[TMP48]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP14]], i32** [[TMP49]], align 8
|
|
// CHECK1-NEXT: [[TMP50:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP50]], align 8
|
|
// CHECK1-NEXT: [[TMP51:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4
|
|
// CHECK1-NEXT: [[TMP52:%.*]] = bitcast i8** [[TMP51]] to i32**
|
|
// CHECK1-NEXT: store i32* [[A]], i32** [[TMP52]], align 8
|
|
// CHECK1-NEXT: [[TMP53:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 4
|
|
// CHECK1-NEXT: [[TMP54:%.*]] = bitcast i8** [[TMP53]] to i32**
|
|
// CHECK1-NEXT: store i32* [[A]], i32** [[TMP54]], align 8
|
|
// CHECK1-NEXT: [[TMP55:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 4
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP55]], align 8
|
|
// CHECK1-NEXT: [[TMP56:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5
|
|
// CHECK1-NEXT: [[TMP57:%.*]] = bitcast i8** [[TMP56]] to %class.anon**
|
|
// CHECK1-NEXT: store %class.anon* [[TMP15]], %class.anon** [[TMP57]], align 8
|
|
// CHECK1-NEXT: [[TMP58:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 5
|
|
// CHECK1-NEXT: [[TMP59:%.*]] = bitcast i8** [[TMP58]] to %class.anon**
|
|
// CHECK1-NEXT: store %class.anon* [[TMP15]], %class.anon** [[TMP59]], align 8
|
|
// CHECK1-NEXT: [[TMP60:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 5
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP60]], align 8
|
|
// CHECK1-NEXT: [[TMP61:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 6
|
|
// CHECK1-NEXT: [[TMP62:%.*]] = bitcast i8** [[TMP61]] to i32***
|
|
// CHECK1-NEXT: store i32** [[TMP16]], i32*** [[TMP62]], align 8
|
|
// CHECK1-NEXT: [[TMP63:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 6
|
|
// CHECK1-NEXT: [[TMP64:%.*]] = bitcast i8** [[TMP63]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP18]], i32** [[TMP64]], align 8
|
|
// CHECK1-NEXT: [[TMP65:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 6
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP65]], align 8
|
|
// CHECK1-NEXT: [[TMP66:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 7
|
|
// CHECK1-NEXT: [[TMP67:%.*]] = bitcast i8** [[TMP66]] to i32***
|
|
// CHECK1-NEXT: store i32** [[TMP19]], i32*** [[TMP67]], align 8
|
|
// CHECK1-NEXT: [[TMP68:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 7
|
|
// CHECK1-NEXT: [[TMP69:%.*]] = bitcast i8** [[TMP68]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP21]], i32** [[TMP69]], align 8
|
|
// CHECK1-NEXT: [[TMP70:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 7
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP70]], align 8
|
|
// CHECK1-NEXT: [[TMP71:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 8
|
|
// CHECK1-NEXT: [[TMP72:%.*]] = bitcast i8** [[TMP71]] to i32***
|
|
// CHECK1-NEXT: store i32** [[TMP22]], i32*** [[TMP72]], align 8
|
|
// CHECK1-NEXT: [[TMP73:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 8
|
|
// CHECK1-NEXT: [[TMP74:%.*]] = bitcast i8** [[TMP73]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP24]], i32** [[TMP74]], align 8
|
|
// CHECK1-NEXT: [[TMP75:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 8
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP75]], align 8
|
|
// CHECK1-NEXT: [[TMP76:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 9
|
|
// CHECK1-NEXT: [[TMP77:%.*]] = bitcast i8** [[TMP76]] to i32****
|
|
// CHECK1-NEXT: store i32*** [[TMP25]], i32**** [[TMP77]], align 8
|
|
// CHECK1-NEXT: [[TMP78:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 9
|
|
// CHECK1-NEXT: [[TMP79:%.*]] = bitcast i8** [[TMP78]] to i32***
|
|
// CHECK1-NEXT: store i32** [[TMP27]], i32*** [[TMP79]], align 8
|
|
// CHECK1-NEXT: [[TMP80:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 9
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP80]], align 8
|
|
// CHECK1-NEXT: [[TMP81:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 10
|
|
// CHECK1-NEXT: [[TMP82:%.*]] = bitcast i8** [[TMP81]] to i32***
|
|
// CHECK1-NEXT: store i32** [[TMP28]], i32*** [[TMP82]], align 8
|
|
// CHECK1-NEXT: [[TMP83:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 10
|
|
// CHECK1-NEXT: [[TMP84:%.*]] = bitcast i8** [[TMP83]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP30]], i32** [[TMP84]], align 8
|
|
// CHECK1-NEXT: [[TMP85:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 10
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP85]], align 8
|
|
// CHECK1-NEXT: [[TMP86:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP87:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP88:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l41.region_id, i32 11, i8** [[TMP86]], i8** [[TMP87]], i64* getelementptr inbounds ([11 x i64], [11 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([11 x i64], [11 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null)
|
|
// CHECK1-NEXT: [[TMP89:%.*]] = icmp ne i32 [[TMP88]], 0
|
|
// CHECK1-NEXT: br i1 [[TMP89]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
|
|
// CHECK1: omp_offload.failed:
|
|
// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l41(i64 [[TMP11]], i32* [[TMP12]], i32* [[TMP13]], i32* [[TMP14]], i32* [[A]], %class.anon* [[TMP15]]) #[[ATTR4:[0-9]+]]
|
|
// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]
|
|
// CHECK1: omp_offload.cont:
|
|
// CHECK1-NEXT: [[TMP90:%.*]] = load i32*, i32** [[B]], align 8
|
|
// CHECK1-NEXT: store i32* [[TMP90]], i32** [[_TMP4]], align 8
|
|
// CHECK1-NEXT: [[TMP91:%.*]] = load i32*, i32** [[C]], align 8
|
|
// CHECK1-NEXT: store i32* [[TMP91]], i32** [[_TMP5]], align 8
|
|
// CHECK1-NEXT: [[TMP92:%.*]] = load %class.anon*, %class.anon** [[L]], align 8
|
|
// CHECK1-NEXT: store %class.anon* [[TMP92]], %class.anon** [[_TMP6]], align 8
|
|
// CHECK1-NEXT: [[TMP93:%.*]] = load i32*, i32** [[_TMP4]], align 8
|
|
// CHECK1-NEXT: [[TMP94:%.*]] = load i32*, i32** [[_TMP5]], align 8
|
|
// CHECK1-NEXT: [[TMP95:%.*]] = load i32*, i32** [[D]], align 8
|
|
// CHECK1-NEXT: [[TMP96:%.*]] = load %class.anon*, %class.anon** [[_TMP6]], align 8
|
|
// CHECK1-NEXT: [[TMP97:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP96]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP98:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP96]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP99:%.*]] = load i32*, i32** [[TMP98]], align 8
|
|
// CHECK1-NEXT: [[TMP100:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP96]], i32 0, i32 1
|
|
// CHECK1-NEXT: [[TMP101:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP96]], i32 0, i32 1
|
|
// CHECK1-NEXT: [[TMP102:%.*]] = load i32*, i32** [[TMP101]], align 8
|
|
// CHECK1-NEXT: [[TMP103:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP96]], i32 0, i32 2
|
|
// CHECK1-NEXT: [[TMP104:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP96]], i32 0, i32 2
|
|
// CHECK1-NEXT: [[TMP105:%.*]] = load i32*, i32** [[TMP104]], align 8
|
|
// CHECK1-NEXT: [[TMP106:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP96]], i32 0, i32 3
|
|
// CHECK1-NEXT: [[TMP107:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP96]], i32 0, i32 3
|
|
// CHECK1-NEXT: [[TMP108:%.*]] = load i32**, i32*** [[TMP107]], align 8
|
|
// CHECK1-NEXT: [[TMP109:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP96]], i32 0, i32 4
|
|
// CHECK1-NEXT: [[TMP110:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP96]], i32 0, i32 4
|
|
// CHECK1-NEXT: [[TMP111:%.*]] = load i32*, i32** [[TMP110]], align 8
|
|
// CHECK1-NEXT: [[TMP112:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP113:%.*]] = bitcast i8** [[TMP112]] to i32**
|
|
// CHECK1-NEXT: store i32* [[ARGC_ADDR]], i32** [[TMP113]], align 8
|
|
// CHECK1-NEXT: [[TMP114:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS8]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP115:%.*]] = bitcast i8** [[TMP114]] to i32**
|
|
// CHECK1-NEXT: store i32* [[ARGC_ADDR]], i32** [[TMP115]], align 8
|
|
// CHECK1-NEXT: [[TMP116:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 0
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP116]], align 8
|
|
// CHECK1-NEXT: [[TMP117:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 1
|
|
// CHECK1-NEXT: [[TMP118:%.*]] = bitcast i8** [[TMP117]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP93]], i32** [[TMP118]], align 8
|
|
// CHECK1-NEXT: [[TMP119:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS8]], i32 0, i32 1
|
|
// CHECK1-NEXT: [[TMP120:%.*]] = bitcast i8** [[TMP119]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP93]], i32** [[TMP120]], align 8
|
|
// CHECK1-NEXT: [[TMP121:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 1
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP121]], align 8
|
|
// CHECK1-NEXT: [[TMP122:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 2
|
|
// CHECK1-NEXT: [[TMP123:%.*]] = bitcast i8** [[TMP122]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP94]], i32** [[TMP123]], align 8
|
|
// CHECK1-NEXT: [[TMP124:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS8]], i32 0, i32 2
|
|
// CHECK1-NEXT: [[TMP125:%.*]] = bitcast i8** [[TMP124]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP94]], i32** [[TMP125]], align 8
|
|
// CHECK1-NEXT: [[TMP126:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 2
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP126]], align 8
|
|
// CHECK1-NEXT: [[TMP127:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 3
|
|
// CHECK1-NEXT: [[TMP128:%.*]] = bitcast i8** [[TMP127]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP95]], i32** [[TMP128]], align 8
|
|
// CHECK1-NEXT: [[TMP129:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS8]], i32 0, i32 3
|
|
// CHECK1-NEXT: [[TMP130:%.*]] = bitcast i8** [[TMP129]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP95]], i32** [[TMP130]], align 8
|
|
// CHECK1-NEXT: [[TMP131:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 3
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP131]], align 8
|
|
// CHECK1-NEXT: [[TMP132:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 4
|
|
// CHECK1-NEXT: [[TMP133:%.*]] = bitcast i8** [[TMP132]] to i32**
|
|
// CHECK1-NEXT: store i32* [[A]], i32** [[TMP133]], align 8
|
|
// CHECK1-NEXT: [[TMP134:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS8]], i32 0, i32 4
|
|
// CHECK1-NEXT: [[TMP135:%.*]] = bitcast i8** [[TMP134]] to i32**
|
|
// CHECK1-NEXT: store i32* [[A]], i32** [[TMP135]], align 8
|
|
// CHECK1-NEXT: [[TMP136:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 4
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP136]], align 8
|
|
// CHECK1-NEXT: [[TMP137:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 5
|
|
// CHECK1-NEXT: [[TMP138:%.*]] = bitcast i8** [[TMP137]] to %class.anon**
|
|
// CHECK1-NEXT: store %class.anon* [[TMP96]], %class.anon** [[TMP138]], align 8
|
|
// CHECK1-NEXT: [[TMP139:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS8]], i32 0, i32 5
|
|
// CHECK1-NEXT: [[TMP140:%.*]] = bitcast i8** [[TMP139]] to %class.anon**
|
|
// CHECK1-NEXT: store %class.anon* [[TMP96]], %class.anon** [[TMP140]], align 8
|
|
// CHECK1-NEXT: [[TMP141:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 5
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP141]], align 8
|
|
// CHECK1-NEXT: [[TMP142:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 6
|
|
// CHECK1-NEXT: [[TMP143:%.*]] = bitcast i8** [[TMP142]] to i32***
|
|
// CHECK1-NEXT: store i32** [[TMP97]], i32*** [[TMP143]], align 8
|
|
// CHECK1-NEXT: [[TMP144:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS8]], i32 0, i32 6
|
|
// CHECK1-NEXT: [[TMP145:%.*]] = bitcast i8** [[TMP144]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP99]], i32** [[TMP145]], align 8
|
|
// CHECK1-NEXT: [[TMP146:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 6
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP146]], align 8
|
|
// CHECK1-NEXT: [[TMP147:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 7
|
|
// CHECK1-NEXT: [[TMP148:%.*]] = bitcast i8** [[TMP147]] to i32***
|
|
// CHECK1-NEXT: store i32** [[TMP100]], i32*** [[TMP148]], align 8
|
|
// CHECK1-NEXT: [[TMP149:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS8]], i32 0, i32 7
|
|
// CHECK1-NEXT: [[TMP150:%.*]] = bitcast i8** [[TMP149]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP102]], i32** [[TMP150]], align 8
|
|
// CHECK1-NEXT: [[TMP151:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 7
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP151]], align 8
|
|
// CHECK1-NEXT: [[TMP152:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 8
|
|
// CHECK1-NEXT: [[TMP153:%.*]] = bitcast i8** [[TMP152]] to i32***
|
|
// CHECK1-NEXT: store i32** [[TMP103]], i32*** [[TMP153]], align 8
|
|
// CHECK1-NEXT: [[TMP154:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS8]], i32 0, i32 8
|
|
// CHECK1-NEXT: [[TMP155:%.*]] = bitcast i8** [[TMP154]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP105]], i32** [[TMP155]], align 8
|
|
// CHECK1-NEXT: [[TMP156:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 8
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP156]], align 8
|
|
// CHECK1-NEXT: [[TMP157:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 9
|
|
// CHECK1-NEXT: [[TMP158:%.*]] = bitcast i8** [[TMP157]] to i32****
|
|
// CHECK1-NEXT: store i32*** [[TMP106]], i32**** [[TMP158]], align 8
|
|
// CHECK1-NEXT: [[TMP159:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS8]], i32 0, i32 9
|
|
// CHECK1-NEXT: [[TMP160:%.*]] = bitcast i8** [[TMP159]] to i32***
|
|
// CHECK1-NEXT: store i32** [[TMP108]], i32*** [[TMP160]], align 8
|
|
// CHECK1-NEXT: [[TMP161:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 9
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP161]], align 8
|
|
// CHECK1-NEXT: [[TMP162:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 10
|
|
// CHECK1-NEXT: [[TMP163:%.*]] = bitcast i8** [[TMP162]] to i32***
|
|
// CHECK1-NEXT: store i32** [[TMP109]], i32*** [[TMP163]], align 8
|
|
// CHECK1-NEXT: [[TMP164:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS8]], i32 0, i32 10
|
|
// CHECK1-NEXT: [[TMP165:%.*]] = bitcast i8** [[TMP164]] to i32**
|
|
// CHECK1-NEXT: store i32* [[TMP111]], i32** [[TMP165]], align 8
|
|
// CHECK1-NEXT: [[TMP166:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS9]], i64 0, i64 10
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP166]], align 8
|
|
// CHECK1-NEXT: [[TMP167:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS7]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP168:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS8]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP169:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l43.region_id, i32 11, i8** [[TMP167]], i8** [[TMP168]], i64* getelementptr inbounds ([11 x i64], [11 x i64]* @.offload_sizes.1, i32 0, i32 0), i64* getelementptr inbounds ([11 x i64], [11 x i64]* @.offload_maptypes.2, i32 0, i32 0), i8** null, i8** null, i32 1, i32 0)
|
|
// CHECK1-NEXT: [[TMP170:%.*]] = icmp ne i32 [[TMP169]], 0
|
|
// CHECK1-NEXT: br i1 [[TMP170]], label [[OMP_OFFLOAD_FAILED10:%.*]], label [[OMP_OFFLOAD_CONT11:%.*]]
|
|
// CHECK1: omp_offload.failed10:
|
|
// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l43(i32* [[ARGC_ADDR]], i32* [[TMP93]], i32* [[TMP94]], i32* [[TMP95]], i32* [[A]], %class.anon* [[TMP96]]) #[[ATTR4]]
|
|
// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT11]]
|
|
// CHECK1: omp_offload.cont11:
|
|
// CHECK1-NEXT: [[TMP171:%.*]] = load i32, i32* [[ARGC_ADDR]], align 4
|
|
// CHECK1-NEXT: [[CALL:%.*]] = call signext i32 @_ZN1S3fooEv(%struct.S* nonnull dereferenceable(4) @s)
|
|
// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP171]], [[CALL]]
|
|
// CHECK1-NEXT: ret i32 [[ADD]]
|
|
//
|
|
//
|
|
// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l41
|
|
// CHECK1-SAME: (i64 [[ARGC:%.*]], i32* nonnull align 4 dereferenceable(4) [[B:%.*]], i32* nonnull align 4 dereferenceable(4) [[C:%.*]], i32* [[D:%.*]], i32* nonnull align 4 dereferenceable(4) [[A:%.*]], %class.anon* nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR1:[0-9]+]] {
|
|
// CHECK1-NEXT: entry:
|
|
// CHECK1-NEXT: [[ARGC_ADDR:%.*]] = alloca i64, align 8
|
|
// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[D_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[L_ADDR:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK1-NEXT: [[TMP:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[_TMP1:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[_TMP2:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK1-NEXT: [[L3:%.*]] = alloca [[CLASS_ANON:%.*]], align 8
|
|
// CHECK1-NEXT: [[_TMP4:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK1-NEXT: [[B5:%.*]] = alloca i32, align 4
|
|
// CHECK1-NEXT: [[_TMP6:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[C7:%.*]] = alloca i32, align 4
|
|
// CHECK1-NEXT: [[_TMP8:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: store i64 [[ARGC]], i64* [[ARGC_ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[B]], i32** [[B_ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[D]], i32** [[D_ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[A]], i32** [[A_ADDR]], align 8
|
|
// CHECK1-NEXT: store %class.anon* [[L]], %class.anon** [[L_ADDR]], align 8
|
|
// CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[ARGC_ADDR]] to i32*
|
|
// CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[B_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP1:%.*]] = load i32*, i32** [[C_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP2:%.*]] = load i32*, i32** [[A_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP3:%.*]] = load %class.anon*, %class.anon** [[L_ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[TMP0]], i32** [[TMP]], align 8
|
|
// CHECK1-NEXT: store i32* [[TMP1]], i32** [[_TMP1]], align 8
|
|
// CHECK1-NEXT: store %class.anon* [[TMP3]], %class.anon** [[_TMP2]], align 8
|
|
// CHECK1-NEXT: [[TMP4:%.*]] = load %class.anon*, %class.anon** [[_TMP2]], align 8
|
|
// CHECK1-NEXT: [[TMP5:%.*]] = bitcast %class.anon* [[L3]] to i8*
|
|
// CHECK1-NEXT: [[TMP6:%.*]] = bitcast %class.anon* [[TMP4]] to i8*
|
|
// CHECK1-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP5]], i8* align 8 [[TMP6]], i64 40, i1 false)
|
|
// CHECK1-NEXT: store %class.anon* [[L3]], %class.anon** [[_TMP4]], align 8
|
|
// CHECK1-NEXT: [[TMP7:%.*]] = load i32*, i32** [[TMP]], align 8
|
|
// CHECK1-NEXT: [[TMP8:%.*]] = load i32, i32* [[TMP7]], align 4
|
|
// CHECK1-NEXT: store i32 [[TMP8]], i32* [[B5]], align 4
|
|
// CHECK1-NEXT: store i32* [[B5]], i32** [[_TMP6]], align 8
|
|
// CHECK1-NEXT: [[TMP9:%.*]] = load i32*, i32** [[_TMP1]], align 8
|
|
// CHECK1-NEXT: [[TMP10:%.*]] = load i32, i32* [[TMP9]], align 4
|
|
// CHECK1-NEXT: store i32 [[TMP10]], i32* [[C7]], align 4
|
|
// CHECK1-NEXT: store i32* [[C7]], i32** [[_TMP8]], align 8
|
|
// CHECK1-NEXT: [[TMP11:%.*]] = load %class.anon*, %class.anon** [[_TMP4]], align 8
|
|
// CHECK1-NEXT: [[CALL:%.*]] = call i64 @"_ZZ4mainENK3$_0clEv"(%class.anon* nonnull dereferenceable(40) [[TMP11]])
|
|
// CHECK1-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l43
|
|
// CHECK1-SAME: (i32* nonnull align 4 dereferenceable(4) [[ARGC:%.*]], i32* nonnull align 4 dereferenceable(4) [[B:%.*]], i32* nonnull align 4 dereferenceable(4) [[C:%.*]], i32* [[D:%.*]], i32* nonnull align 4 dereferenceable(4) [[A:%.*]], %class.anon* nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK1-NEXT: entry:
|
|
// CHECK1-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[D_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[L_ADDR:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK1-NEXT: [[TMP:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[_TMP1:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[_TMP2:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK1-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[B]], i32** [[B_ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[D]], i32** [[D_ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[A]], i32** [[A_ADDR]], align 8
|
|
// CHECK1-NEXT: store %class.anon* [[L]], %class.anon** [[L_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP1:%.*]] = load i32*, i32** [[B_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP2:%.*]] = load i32*, i32** [[C_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP3:%.*]] = load i32*, i32** [[A_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP4:%.*]] = load %class.anon*, %class.anon** [[L_ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[TMP1]], i32** [[TMP]], align 8
|
|
// CHECK1-NEXT: store i32* [[TMP2]], i32** [[_TMP1]], align 8
|
|
// CHECK1-NEXT: store %class.anon* [[TMP4]], %class.anon** [[_TMP2]], align 8
|
|
// CHECK1-NEXT: [[TMP5:%.*]] = load i32*, i32** [[TMP]], align 8
|
|
// CHECK1-NEXT: [[TMP6:%.*]] = load i32*, i32** [[_TMP1]], align 8
|
|
// CHECK1-NEXT: [[TMP7:%.*]] = load i32*, i32** [[D_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP8:%.*]] = load %class.anon*, %class.anon** [[_TMP2]], align 8
|
|
// CHECK1-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1]], i32 6, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*, i32*, i32*, i32*, i32*, %class.anon*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[TMP0]], i32* [[TMP5]], i32* [[TMP6]], i32* [[TMP7]], i32* [[TMP3]], %class.anon* [[TMP8]])
|
|
// CHECK1-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK1-LABEL: define {{[^@]+}}@.omp_outlined.
|
|
// CHECK1-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[ARGC:%.*]], i32* nonnull align 4 dereferenceable(4) [[B:%.*]], i32* nonnull align 4 dereferenceable(4) [[C:%.*]], i32* [[D:%.*]], i32* nonnull align 4 dereferenceable(4) [[A:%.*]], %class.anon* nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK1-NEXT: entry:
|
|
// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[D_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[L_ADDR:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK1-NEXT: [[TMP:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[_TMP1:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[_TMP2:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK1-NEXT: [[L3:%.*]] = alloca [[CLASS_ANON:%.*]], align 8
|
|
// CHECK1-NEXT: [[_TMP4:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK1-NEXT: [[ARGC5:%.*]] = alloca i32, align 4
|
|
// CHECK1-NEXT: [[B6:%.*]] = alloca i32, align 4
|
|
// CHECK1-NEXT: [[_TMP7:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[C8:%.*]] = alloca i32, align 4
|
|
// CHECK1-NEXT: [[_TMP9:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[A10:%.*]] = alloca i32, align 4
|
|
// CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[B]], i32** [[B_ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[D]], i32** [[D_ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[A]], i32** [[A_ADDR]], align 8
|
|
// CHECK1-NEXT: store %class.anon* [[L]], %class.anon** [[L_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP1:%.*]] = load i32*, i32** [[B_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP2:%.*]] = load i32*, i32** [[C_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP3:%.*]] = load i32*, i32** [[A_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP4:%.*]] = load %class.anon*, %class.anon** [[L_ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[TMP1]], i32** [[TMP]], align 8
|
|
// CHECK1-NEXT: store i32* [[TMP2]], i32** [[_TMP1]], align 8
|
|
// CHECK1-NEXT: store %class.anon* [[TMP4]], %class.anon** [[_TMP2]], align 8
|
|
// CHECK1-NEXT: [[TMP5:%.*]] = load %class.anon*, %class.anon** [[_TMP2]], align 8
|
|
// CHECK1-NEXT: [[TMP6:%.*]] = bitcast %class.anon* [[L3]] to i8*
|
|
// CHECK1-NEXT: [[TMP7:%.*]] = bitcast %class.anon* [[TMP5]] to i8*
|
|
// CHECK1-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP6]], i8* align 8 [[TMP7]], i64 40, i1 false)
|
|
// CHECK1-NEXT: store %class.anon* [[L3]], %class.anon** [[_TMP4]], align 8
|
|
// CHECK1-NEXT: [[TMP8:%.*]] = load i32, i32* [[TMP0]], align 4
|
|
// CHECK1-NEXT: store i32 [[TMP8]], i32* [[ARGC5]], align 4
|
|
// CHECK1-NEXT: [[TMP9:%.*]] = load i32*, i32** [[TMP]], align 8
|
|
// CHECK1-NEXT: [[TMP10:%.*]] = load i32, i32* [[TMP9]], align 4
|
|
// CHECK1-NEXT: store i32 [[TMP10]], i32* [[B6]], align 4
|
|
// CHECK1-NEXT: store i32* [[B6]], i32** [[_TMP7]], align 8
|
|
// CHECK1-NEXT: [[TMP11:%.*]] = load i32*, i32** [[_TMP1]], align 8
|
|
// CHECK1-NEXT: [[TMP12:%.*]] = load i32, i32* [[TMP11]], align 4
|
|
// CHECK1-NEXT: store i32 [[TMP12]], i32* [[C8]], align 4
|
|
// CHECK1-NEXT: store i32* [[C8]], i32** [[_TMP9]], align 8
|
|
// CHECK1-NEXT: [[TMP13:%.*]] = load i32, i32* [[TMP3]], align 4
|
|
// CHECK1-NEXT: store i32 [[TMP13]], i32* [[A10]], align 4
|
|
// CHECK1-NEXT: [[TMP14:%.*]] = load %class.anon*, %class.anon** [[_TMP4]], align 8
|
|
// CHECK1-NEXT: [[CALL:%.*]] = call i64 @"_ZZ4mainENK3$_0clEv"(%class.anon* nonnull dereferenceable(40) [[TMP14]])
|
|
// CHECK1-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK1-LABEL: define {{[^@]+}}@_ZN1S3fooEv
|
|
// CHECK1-SAME: (%struct.S* nonnull dereferenceable(4) [[THIS:%.*]]) #[[ATTR3:[0-9]+]] comdat align 2 {
|
|
// CHECK1-NEXT: entry:
|
|
// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.S*, align 8
|
|
// CHECK1-NEXT: [[L:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8
|
|
// CHECK1-NEXT: [[TMP:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x i8*], align 8
|
|
// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x i8*], align 8
|
|
// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x i8*], align 8
|
|
// CHECK1-NEXT: [[_TMP2:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [3 x i8*], align 8
|
|
// CHECK1-NEXT: [[DOTOFFLOAD_PTRS4:%.*]] = alloca [3 x i8*], align 8
|
|
// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [3 x i8*], align 8
|
|
// CHECK1-NEXT: store %struct.S* [[THIS]], %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK1-NEXT: [[THIS1:%.*]] = load %struct.S*, %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP0:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[REF_TMP]], i32 0, i32 0
|
|
// CHECK1-NEXT: store %struct.S* [[THIS1]], %struct.S** [[TMP0]], align 8
|
|
// CHECK1-NEXT: store %class.anon.0* [[REF_TMP]], %class.anon.0** [[L]], align 8
|
|
// CHECK1-NEXT: [[TMP1:%.*]] = load %class.anon.0*, %class.anon.0** [[L]], align 8
|
|
// CHECK1-NEXT: store %class.anon.0* [[TMP1]], %class.anon.0** [[TMP]], align 8
|
|
// CHECK1-NEXT: [[TMP2:%.*]] = load %class.anon.0*, %class.anon.0** [[TMP]], align 8
|
|
// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP2]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP2]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP6:%.*]] = bitcast i8** [[TMP5]] to %struct.S**
|
|
// CHECK1-NEXT: store %struct.S* [[THIS1]], %struct.S** [[TMP6]], align 8
|
|
// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP8:%.*]] = bitcast i8** [[TMP7]] to %struct.S**
|
|
// CHECK1-NEXT: store %struct.S* [[THIS1]], %struct.S** [[TMP8]], align 8
|
|
// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP9]], align 8
|
|
// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
|
|
// CHECK1-NEXT: [[TMP11:%.*]] = bitcast i8** [[TMP10]] to %class.anon.0**
|
|
// CHECK1-NEXT: store %class.anon.0* [[TMP2]], %class.anon.0** [[TMP11]], align 8
|
|
// CHECK1-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1
|
|
// CHECK1-NEXT: [[TMP13:%.*]] = bitcast i8** [[TMP12]] to %class.anon.0**
|
|
// CHECK1-NEXT: store %class.anon.0* [[TMP2]], %class.anon.0** [[TMP13]], align 8
|
|
// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP14]], align 8
|
|
// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2
|
|
// CHECK1-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to %struct.S***
|
|
// CHECK1-NEXT: store %struct.S** [[TMP3]], %struct.S*** [[TMP16]], align 8
|
|
// CHECK1-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2
|
|
// CHECK1-NEXT: [[TMP18:%.*]] = bitcast i8** [[TMP17]] to %struct.S***
|
|
// CHECK1-NEXT: store %struct.S** [[TMP4]], %struct.S*** [[TMP18]], align 8
|
|
// CHECK1-NEXT: [[TMP19:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP19]], align 8
|
|
// CHECK1-NEXT: [[TMP20:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP21:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP22:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l27.region_id, i32 3, i8** [[TMP20]], i8** [[TMP21]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_sizes.3, i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_maptypes.4, i32 0, i32 0), i8** null, i8** null)
|
|
// CHECK1-NEXT: [[TMP23:%.*]] = icmp ne i32 [[TMP22]], 0
|
|
// CHECK1-NEXT: br i1 [[TMP23]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
|
|
// CHECK1: omp_offload.failed:
|
|
// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l27(%struct.S* [[THIS1]], %class.anon.0* [[TMP2]]) #[[ATTR4]]
|
|
// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]
|
|
// CHECK1: omp_offload.cont:
|
|
// CHECK1-NEXT: [[TMP24:%.*]] = load %class.anon.0*, %class.anon.0** [[L]], align 8
|
|
// CHECK1-NEXT: store %class.anon.0* [[TMP24]], %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK1-NEXT: [[TMP25:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK1-NEXT: [[TMP26:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP25]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP27:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP25]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP28:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP29:%.*]] = bitcast i8** [[TMP28]] to %struct.S**
|
|
// CHECK1-NEXT: store %struct.S* [[THIS1]], %struct.S** [[TMP29]], align 8
|
|
// CHECK1-NEXT: [[TMP30:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP31:%.*]] = bitcast i8** [[TMP30]] to %struct.S**
|
|
// CHECK1-NEXT: store %struct.S* [[THIS1]], %struct.S** [[TMP31]], align 8
|
|
// CHECK1-NEXT: [[TMP32:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS5]], i64 0, i64 0
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP32]], align 8
|
|
// CHECK1-NEXT: [[TMP33:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 1
|
|
// CHECK1-NEXT: [[TMP34:%.*]] = bitcast i8** [[TMP33]] to %class.anon.0**
|
|
// CHECK1-NEXT: store %class.anon.0* [[TMP25]], %class.anon.0** [[TMP34]], align 8
|
|
// CHECK1-NEXT: [[TMP35:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 1
|
|
// CHECK1-NEXT: [[TMP36:%.*]] = bitcast i8** [[TMP35]] to %class.anon.0**
|
|
// CHECK1-NEXT: store %class.anon.0* [[TMP25]], %class.anon.0** [[TMP36]], align 8
|
|
// CHECK1-NEXT: [[TMP37:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS5]], i64 0, i64 1
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP37]], align 8
|
|
// CHECK1-NEXT: [[TMP38:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 2
|
|
// CHECK1-NEXT: [[TMP39:%.*]] = bitcast i8** [[TMP38]] to %struct.S***
|
|
// CHECK1-NEXT: store %struct.S** [[TMP26]], %struct.S*** [[TMP39]], align 8
|
|
// CHECK1-NEXT: [[TMP40:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 2
|
|
// CHECK1-NEXT: [[TMP41:%.*]] = bitcast i8** [[TMP40]] to %struct.S***
|
|
// CHECK1-NEXT: store %struct.S** [[TMP27]], %struct.S*** [[TMP41]], align 8
|
|
// CHECK1-NEXT: [[TMP42:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS5]], i64 0, i64 2
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP42]], align 8
|
|
// CHECK1-NEXT: [[TMP43:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP44:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP45:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l29.region_id, i32 3, i8** [[TMP43]], i8** [[TMP44]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_sizes.6, i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_maptypes.7, i32 0, i32 0), i8** null, i8** null, i32 1, i32 0)
|
|
// CHECK1-NEXT: [[TMP46:%.*]] = icmp ne i32 [[TMP45]], 0
|
|
// CHECK1-NEXT: br i1 [[TMP46]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]]
|
|
// CHECK1: omp_offload.failed6:
|
|
// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l29(%struct.S* [[THIS1]], %class.anon.0* [[TMP25]]) #[[ATTR4]]
|
|
// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT7]]
|
|
// CHECK1: omp_offload.cont7:
|
|
// CHECK1-NEXT: [[A:%.*]] = getelementptr inbounds [[STRUCT_S:%.*]], %struct.S* [[THIS1]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP47:%.*]] = load i32, i32* [[A]], align 4
|
|
// CHECK1-NEXT: [[TMP48:%.*]] = load %class.anon.0*, %class.anon.0** [[L]], align 8
|
|
// CHECK1-NEXT: [[CALL:%.*]] = call signext i32 @_Z3fooIZN1S3fooEvEUlvE_EiRKT_(%class.anon.0* nonnull align 8 dereferenceable(8) [[TMP48]])
|
|
// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP47]], [[CALL]]
|
|
// CHECK1-NEXT: ret i32 [[ADD]]
|
|
//
|
|
//
|
|
// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l27
|
|
// CHECK1-SAME: (%struct.S* [[THIS:%.*]], %class.anon.0* nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK1-NEXT: entry:
|
|
// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.S*, align 8
|
|
// CHECK1-NEXT: [[L_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: [[TMP:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: [[L1:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8
|
|
// CHECK1-NEXT: [[_TMP2:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: store %struct.S* [[THIS]], %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK1-NEXT: store %class.anon.0* [[L]], %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP0:%.*]] = load %struct.S*, %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP1:%.*]] = load %class.anon.0*, %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK1-NEXT: store %class.anon.0* [[TMP1]], %class.anon.0** [[TMP]], align 8
|
|
// CHECK1-NEXT: [[TMP2:%.*]] = load %class.anon.0*, %class.anon.0** [[TMP]], align 8
|
|
// CHECK1-NEXT: [[TMP3:%.*]] = bitcast %class.anon.0* [[L1]] to i8*
|
|
// CHECK1-NEXT: [[TMP4:%.*]] = bitcast %class.anon.0* [[TMP2]] to i8*
|
|
// CHECK1-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP3]], i8* align 8 [[TMP4]], i64 8, i1 false)
|
|
// CHECK1-NEXT: store %class.anon.0* [[L1]], %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK1-NEXT: [[TMP5:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK1-NEXT: [[CALL:%.*]] = call signext i32 @_ZZN1S3fooEvENKUlvE_clEv(%class.anon.0* nonnull dereferenceable(8) [[TMP5]])
|
|
// CHECK1-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK1-LABEL: define {{[^@]+}}@_ZZN1S3fooEvENKUlvE_clEv
|
|
// CHECK1-SAME: (%class.anon.0* nonnull dereferenceable(8) [[THIS:%.*]]) #[[ATTR3]] comdat align 2 {
|
|
// CHECK1-NEXT: entry:
|
|
// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: store %class.anon.0* [[THIS]], %class.anon.0** [[THIS_ADDR]], align 8
|
|
// CHECK1-NEXT: [[THIS1:%.*]] = load %class.anon.0*, %class.anon.0** [[THIS_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP0:%.*]] = getelementptr inbounds [[CLASS_ANON_0:%.*]], %class.anon.0* [[THIS1]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP1:%.*]] = load %struct.S*, %struct.S** [[TMP0]], align 8
|
|
// CHECK1-NEXT: [[A:%.*]] = getelementptr inbounds [[STRUCT_S:%.*]], %struct.S* [[TMP1]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP2:%.*]] = load i32, i32* [[A]], align 4
|
|
// CHECK1-NEXT: ret i32 [[TMP2]]
|
|
//
|
|
//
|
|
// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l29
|
|
// CHECK1-SAME: (%struct.S* [[THIS:%.*]], %class.anon.0* nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK1-NEXT: entry:
|
|
// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.S*, align 8
|
|
// CHECK1-NEXT: [[L_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: [[TMP:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: store %struct.S* [[THIS]], %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK1-NEXT: store %class.anon.0* [[L]], %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP0:%.*]] = load %struct.S*, %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP1:%.*]] = load %class.anon.0*, %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK1-NEXT: store %class.anon.0* [[TMP1]], %class.anon.0** [[TMP]], align 8
|
|
// CHECK1-NEXT: [[TMP2:%.*]] = load %class.anon.0*, %class.anon.0** [[TMP]], align 8
|
|
// CHECK1-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1]], i32 2, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, %struct.S*, %class.anon.0*)* @.omp_outlined..5 to void (i32*, i32*, ...)*), %struct.S* [[TMP0]], %class.anon.0* [[TMP2]])
|
|
// CHECK1-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK1-LABEL: define {{[^@]+}}@.omp_outlined..5
|
|
// CHECK1-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], %struct.S* [[THIS:%.*]], %class.anon.0* nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK1-NEXT: entry:
|
|
// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.S*, align 8
|
|
// CHECK1-NEXT: [[L_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: [[TMP:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: [[L1:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8
|
|
// CHECK1-NEXT: [[_TMP2:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
|
|
// CHECK1-NEXT: store %struct.S* [[THIS]], %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK1-NEXT: store %class.anon.0* [[L]], %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP0:%.*]] = load %struct.S*, %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP1:%.*]] = load %class.anon.0*, %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK1-NEXT: store %class.anon.0* [[TMP1]], %class.anon.0** [[TMP]], align 8
|
|
// CHECK1-NEXT: [[TMP2:%.*]] = load %class.anon.0*, %class.anon.0** [[TMP]], align 8
|
|
// CHECK1-NEXT: [[TMP3:%.*]] = bitcast %class.anon.0* [[L1]] to i8*
|
|
// CHECK1-NEXT: [[TMP4:%.*]] = bitcast %class.anon.0* [[TMP2]] to i8*
|
|
// CHECK1-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP3]], i8* align 8 [[TMP4]], i64 8, i1 false)
|
|
// CHECK1-NEXT: store %class.anon.0* [[L1]], %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK1-NEXT: [[TMP5:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK1-NEXT: [[CALL:%.*]] = call signext i32 @_ZZN1S3fooEvENKUlvE_clEv(%class.anon.0* nonnull dereferenceable(8) [[TMP5]])
|
|
// CHECK1-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK1-LABEL: define {{[^@]+}}@_Z3fooIZN1S3fooEvEUlvE_EiRKT_
|
|
// CHECK1-SAME: (%class.anon.0* nonnull align 8 dereferenceable(8) [[T:%.*]]) #[[ATTR3]] comdat {
|
|
// CHECK1-NEXT: entry:
|
|
// CHECK1-NEXT: [[T_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: [[TMP:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x i8*], align 8
|
|
// CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x i8*], align 8
|
|
// CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x i8*], align 8
|
|
// CHECK1-NEXT: store %class.anon.0* [[T]], %class.anon.0** [[T_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP0:%.*]] = load %class.anon.0*, %class.anon.0** [[T_ADDR]], align 8
|
|
// CHECK1-NEXT: store %class.anon.0* [[TMP0]], %class.anon.0** [[TMP]], align 8
|
|
// CHECK1-NEXT: [[TMP1:%.*]] = load %class.anon.0*, %class.anon.0** [[TMP]], align 8
|
|
// CHECK1-NEXT: [[TMP2:%.*]] = getelementptr inbounds [[CLASS_ANON_0:%.*]], %class.anon.0* [[TMP1]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP1]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP5:%.*]] = bitcast i8** [[TMP4]] to %class.anon.0**
|
|
// CHECK1-NEXT: store %class.anon.0* [[TMP1]], %class.anon.0** [[TMP5]], align 8
|
|
// CHECK1-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP7:%.*]] = bitcast i8** [[TMP6]] to %class.anon.0**
|
|
// CHECK1-NEXT: store %class.anon.0* [[TMP1]], %class.anon.0** [[TMP7]], align 8
|
|
// CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP8]], align 8
|
|
// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
|
|
// CHECK1-NEXT: [[TMP10:%.*]] = bitcast i8** [[TMP9]] to %struct.S***
|
|
// CHECK1-NEXT: store %struct.S** [[TMP2]], %struct.S*** [[TMP10]], align 8
|
|
// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1
|
|
// CHECK1-NEXT: [[TMP12:%.*]] = bitcast i8** [[TMP11]] to %struct.S***
|
|
// CHECK1-NEXT: store %struct.S** [[TMP3]], %struct.S*** [[TMP12]], align 8
|
|
// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
|
|
// CHECK1-NEXT: store i8* null, i8** [[TMP13]], align 8
|
|
// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP15:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
|
|
// CHECK1-NEXT: [[TMP16:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l18.region_id, i32 2, i8** [[TMP14]], i8** [[TMP15]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.9, i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.10, i32 0, i32 0), i8** null, i8** null, i32 1, i32 0)
|
|
// CHECK1-NEXT: [[TMP17:%.*]] = icmp ne i32 [[TMP16]], 0
|
|
// CHECK1-NEXT: br i1 [[TMP17]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
|
|
// CHECK1: omp_offload.failed:
|
|
// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l18(%class.anon.0* [[TMP1]]) #[[ATTR4]]
|
|
// CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]]
|
|
// CHECK1: omp_offload.cont:
|
|
// CHECK1-NEXT: ret i32 0
|
|
//
|
|
//
|
|
// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l18
|
|
// CHECK1-SAME: (%class.anon.0* nonnull align 8 dereferenceable(8) [[T:%.*]]) #[[ATTR1]] {
|
|
// CHECK1-NEXT: entry:
|
|
// CHECK1-NEXT: [[T_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: [[TMP:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: store %class.anon.0* [[T]], %class.anon.0** [[T_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP0:%.*]] = load %class.anon.0*, %class.anon.0** [[T_ADDR]], align 8
|
|
// CHECK1-NEXT: store %class.anon.0* [[TMP0]], %class.anon.0** [[TMP]], align 8
|
|
// CHECK1-NEXT: [[TMP1:%.*]] = load %class.anon.0*, %class.anon.0** [[TMP]], align 8
|
|
// CHECK1-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, %class.anon.0*)* @.omp_outlined..8 to void (i32*, i32*, ...)*), %class.anon.0* [[TMP1]])
|
|
// CHECK1-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK1-LABEL: define {{[^@]+}}@.omp_outlined..8
|
|
// CHECK1-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], %class.anon.0* nonnull align 8 dereferenceable(8) [[T:%.*]]) #[[ATTR1]] {
|
|
// CHECK1-NEXT: entry:
|
|
// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK1-NEXT: [[T_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: [[TMP:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: [[T1:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8
|
|
// CHECK1-NEXT: [[_TMP2:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
|
|
// CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
|
|
// CHECK1-NEXT: store %class.anon.0* [[T]], %class.anon.0** [[T_ADDR]], align 8
|
|
// CHECK1-NEXT: [[TMP0:%.*]] = load %class.anon.0*, %class.anon.0** [[T_ADDR]], align 8
|
|
// CHECK1-NEXT: store %class.anon.0* [[TMP0]], %class.anon.0** [[TMP]], align 8
|
|
// CHECK1-NEXT: [[TMP1:%.*]] = load %class.anon.0*, %class.anon.0** [[TMP]], align 8
|
|
// CHECK1-NEXT: [[TMP2:%.*]] = bitcast %class.anon.0* [[T1]] to i8*
|
|
// CHECK1-NEXT: [[TMP3:%.*]] = bitcast %class.anon.0* [[TMP1]] to i8*
|
|
// CHECK1-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP2]], i8* align 8 [[TMP3]], i64 8, i1 false)
|
|
// CHECK1-NEXT: store %class.anon.0* [[T1]], %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK1-NEXT: [[TMP4:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK1-NEXT: [[CALL:%.*]] = call signext i32 @_ZZN1S3fooEvENKUlvE_clEv(%class.anon.0* nonnull dereferenceable(8) [[TMP4]])
|
|
// CHECK1-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK1-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg
|
|
// CHECK1-SAME: () #[[ATTR5:[0-9]+]] {
|
|
// CHECK1-NEXT: entry:
|
|
// CHECK1-NEXT: call void @__tgt_register_requires(i64 1)
|
|
// CHECK1-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l27_worker
|
|
// CHECK2-SAME: () #[[ATTR0:[0-9]+]] {
|
|
// CHECK2-NEXT: entry:
|
|
// CHECK2-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 8
|
|
// CHECK2-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1
|
|
// CHECK2-NEXT: store i8* null, i8** [[WORK_FN]], align 8
|
|
// CHECK2-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1
|
|
// CHECK2-NEXT: br label [[DOTAWAIT_WORK:%.*]]
|
|
// CHECK2: .await.work:
|
|
// CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
|
|
// CHECK2-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]])
|
|
// CHECK2-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8
|
|
// CHECK2-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1
|
|
// CHECK2-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 8
|
|
// CHECK2-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null
|
|
// CHECK2-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]]
|
|
// CHECK2: .select.workers:
|
|
// CHECK2-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1
|
|
// CHECK2-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0
|
|
// CHECK2-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]]
|
|
// CHECK2: .execute.parallel:
|
|
// CHECK2-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]])
|
|
// CHECK2-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)*
|
|
// CHECK2-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]])
|
|
// CHECK2-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]]
|
|
// CHECK2: .terminate.parallel:
|
|
// CHECK2-NEXT: call void @__kmpc_kernel_end_parallel()
|
|
// CHECK2-NEXT: br label [[DOTBARRIER_PARALLEL]]
|
|
// CHECK2: .barrier.parallel:
|
|
// CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
|
|
// CHECK2-NEXT: br label [[DOTAWAIT_WORK]]
|
|
// CHECK2: .exit:
|
|
// CHECK2-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l27
|
|
// CHECK2-SAME: (%struct.S* [[THIS:%.*]], %class.anon* nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR1:[0-9]+]] {
|
|
// CHECK2-NEXT: entry:
|
|
// CHECK2-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.S*, align 8
|
|
// CHECK2-NEXT: [[L_ADDR:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK2-NEXT: [[TMP:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK2-NEXT: [[L7:%.*]] = alloca [[CLASS_ANON:%.*]], align 8
|
|
// CHECK2-NEXT: [[_TMP8:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK2-NEXT: store %struct.S* [[THIS]], %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK2-NEXT: store %class.anon* [[L]], %class.anon** [[L_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP0:%.*]] = load %struct.S*, %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP1:%.*]] = load %class.anon*, %class.anon** [[L_ADDR]], align 8
|
|
// CHECK2-NEXT: store %class.anon* [[TMP1]], %class.anon** [[TMP]], align 8
|
|
// CHECK2-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
|
|
// CHECK2-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK2-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
|
|
// CHECK2-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]]
|
|
// CHECK2-NEXT: [[TMP2:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]]
|
|
// CHECK2-NEXT: br i1 [[TMP2]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]]
|
|
// CHECK2: .worker:
|
|
// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l27_worker() #[[ATTR6:[0-9]+]]
|
|
// CHECK2-NEXT: br label [[DOTEXIT:%.*]]
|
|
// CHECK2: .mastercheck:
|
|
// CHECK2-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
|
|
// CHECK2-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK2-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
|
|
// CHECK2-NEXT: [[TMP3:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1
|
|
// CHECK2-NEXT: [[TMP4:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1
|
|
// CHECK2-NEXT: [[TMP5:%.*]] = xor i32 [[TMP3]], -1
|
|
// CHECK2-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP4]], [[TMP5]]
|
|
// CHECK2-NEXT: [[TMP6:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]]
|
|
// CHECK2-NEXT: br i1 [[TMP6]], label [[DOTMASTER:%.*]], label [[DOTEXIT]]
|
|
// CHECK2: .master:
|
|
// CHECK2-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK2-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
|
|
// CHECK2-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]]
|
|
// CHECK2-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1)
|
|
// CHECK2-NEXT: call void @__kmpc_data_sharing_init_stack()
|
|
// CHECK2-NEXT: [[TMP7:%.*]] = load %class.anon*, %class.anon** [[TMP]], align 8
|
|
// CHECK2-NEXT: [[TMP8:%.*]] = bitcast %class.anon* [[L7]] to i8*
|
|
// CHECK2-NEXT: [[TMP9:%.*]] = bitcast %class.anon* [[TMP7]] to i8*
|
|
// CHECK2-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP8]], i8* align 8 [[TMP9]], i64 8, i1 false)
|
|
// CHECK2-NEXT: store %class.anon* [[L7]], %class.anon** [[_TMP8]], align 8
|
|
// CHECK2-NEXT: [[TMP10:%.*]] = load %class.anon*, %class.anon** [[_TMP8]], align 8
|
|
// CHECK2-NEXT: [[TMP11:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP10]], i32 0, i32 0
|
|
// CHECK2-NEXT: store %struct.S* [[TMP0]], %struct.S** [[TMP11]], align 8
|
|
// CHECK2-NEXT: [[TMP12:%.*]] = load %class.anon*, %class.anon** [[_TMP8]], align 8
|
|
// CHECK2-NEXT: [[CALL:%.*]] = call i32 @_ZZN1S3fooEvENKUlvE_clEv(%class.anon* nonnull dereferenceable(8) [[TMP12]]) #[[ATTR7:[0-9]+]]
|
|
// CHECK2-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]]
|
|
// CHECK2: .termination.notifier:
|
|
// CHECK2-NEXT: call void @__kmpc_kernel_deinit(i16 1)
|
|
// CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
|
|
// CHECK2-NEXT: br label [[DOTEXIT]]
|
|
// CHECK2: .exit:
|
|
// CHECK2-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK2-LABEL: define {{[^@]+}}@_ZZN1S3fooEvENKUlvE_clEv
|
|
// CHECK2-SAME: (%class.anon* nonnull dereferenceable(8) [[THIS:%.*]]) #[[ATTR4:[0-9]+]] comdat align 2 {
|
|
// CHECK2-NEXT: entry:
|
|
// CHECK2-NEXT: [[THIS_ADDR:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK2-NEXT: store %class.anon* [[THIS]], %class.anon** [[THIS_ADDR]], align 8
|
|
// CHECK2-NEXT: [[THIS1:%.*]] = load %class.anon*, %class.anon** [[THIS_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP0:%.*]] = getelementptr inbounds [[CLASS_ANON:%.*]], %class.anon* [[THIS1]], i32 0, i32 0
|
|
// CHECK2-NEXT: [[TMP1:%.*]] = load %struct.S*, %struct.S** [[TMP0]], align 8
|
|
// CHECK2-NEXT: [[A:%.*]] = getelementptr inbounds [[STRUCT_S:%.*]], %struct.S* [[TMP1]], i32 0, i32 0
|
|
// CHECK2-NEXT: [[TMP2:%.*]] = load i32, i32* [[A]], align 4
|
|
// CHECK2-NEXT: ret i32 [[TMP2]]
|
|
//
|
|
//
|
|
// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l29
|
|
// CHECK2-SAME: (%struct.S* [[THIS:%.*]], %class.anon* nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK2-NEXT: entry:
|
|
// CHECK2-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.S*, align 8
|
|
// CHECK2-NEXT: [[L_ADDR:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK2-NEXT: [[TMP:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [2 x i8*], align 8
|
|
// CHECK2-NEXT: store %struct.S* [[THIS]], %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK2-NEXT: store %class.anon* [[L]], %class.anon** [[L_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP0:%.*]] = load %struct.S*, %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP1:%.*]] = load %class.anon*, %class.anon** [[L_ADDR]], align 8
|
|
// CHECK2-NEXT: store %class.anon* [[TMP1]], %class.anon** [[TMP]], align 8
|
|
// CHECK2-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK2-NEXT: call void @__kmpc_spmd_kernel_init(i32 [[NVPTX_NUM_THREADS]], i16 1)
|
|
// CHECK2-NEXT: call void @__kmpc_data_sharing_init_stack_spmd()
|
|
// CHECK2-NEXT: br label [[DOTEXECUTE:%.*]]
|
|
// CHECK2: .execute:
|
|
// CHECK2-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2:[0-9]+]])
|
|
// CHECK2-NEXT: [[TMP3:%.*]] = load %class.anon*, %class.anon** [[TMP]], align 8
|
|
// CHECK2-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0
|
|
// CHECK2-NEXT: [[TMP5:%.*]] = bitcast %struct.S* [[TMP0]] to i8*
|
|
// CHECK2-NEXT: store i8* [[TMP5]], i8** [[TMP4]], align 8
|
|
// CHECK2-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 1
|
|
// CHECK2-NEXT: [[TMP7:%.*]] = bitcast %class.anon* [[TMP3]] to i8*
|
|
// CHECK2-NEXT: store i8* [[TMP7]], i8** [[TMP6]], align 8
|
|
// CHECK2-NEXT: [[TMP8:%.*]] = bitcast [2 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
|
|
// CHECK2-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP2]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, %struct.S*, %class.anon*)* @__omp_outlined__ to i8*), i8* null, i8** [[TMP8]], i64 2)
|
|
// CHECK2-NEXT: br label [[DOTOMP_DEINIT:%.*]]
|
|
// CHECK2: .omp.deinit:
|
|
// CHECK2-NEXT: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
|
|
// CHECK2-NEXT: br label [[DOTEXIT:%.*]]
|
|
// CHECK2: .exit:
|
|
// CHECK2-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__
|
|
// CHECK2-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], %struct.S* [[THIS:%.*]], %class.anon* nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK2-NEXT: entry:
|
|
// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.S*, align 8
|
|
// CHECK2-NEXT: [[L_ADDR:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK2-NEXT: [[TMP:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK2-NEXT: [[L1:%.*]] = alloca [[CLASS_ANON:%.*]], align 8
|
|
// CHECK2-NEXT: [[_TMP2:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
|
|
// CHECK2-NEXT: store %struct.S* [[THIS]], %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK2-NEXT: store %class.anon* [[L]], %class.anon** [[L_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP0:%.*]] = load %struct.S*, %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP1:%.*]] = load %class.anon*, %class.anon** [[L_ADDR]], align 8
|
|
// CHECK2-NEXT: store %class.anon* [[TMP1]], %class.anon** [[TMP]], align 8
|
|
// CHECK2-NEXT: [[TMP2:%.*]] = load %class.anon*, %class.anon** [[TMP]], align 8
|
|
// CHECK2-NEXT: [[TMP3:%.*]] = bitcast %class.anon* [[L1]] to i8*
|
|
// CHECK2-NEXT: [[TMP4:%.*]] = bitcast %class.anon* [[TMP2]] to i8*
|
|
// CHECK2-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP3]], i8* align 8 [[TMP4]], i64 8, i1 false)
|
|
// CHECK2-NEXT: store %class.anon* [[L1]], %class.anon** [[_TMP2]], align 8
|
|
// CHECK2-NEXT: [[TMP5:%.*]] = load %class.anon*, %class.anon** [[_TMP2]], align 8
|
|
// CHECK2-NEXT: [[TMP6:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP5]], i32 0, i32 0
|
|
// CHECK2-NEXT: store %struct.S* [[TMP0]], %struct.S** [[TMP6]], align 8
|
|
// CHECK2-NEXT: [[TMP7:%.*]] = load %class.anon*, %class.anon** [[_TMP2]], align 8
|
|
// CHECK2-NEXT: [[CALL:%.*]] = call i32 @_ZZN1S3fooEvENKUlvE_clEv(%class.anon* nonnull dereferenceable(8) [[TMP7]]) #[[ATTR7]]
|
|
// CHECK2-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l41_worker
|
|
// CHECK2-SAME: () #[[ATTR0]] {
|
|
// CHECK2-NEXT: entry:
|
|
// CHECK2-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 8
|
|
// CHECK2-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1
|
|
// CHECK2-NEXT: store i8* null, i8** [[WORK_FN]], align 8
|
|
// CHECK2-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1
|
|
// CHECK2-NEXT: br label [[DOTAWAIT_WORK:%.*]]
|
|
// CHECK2: .await.work:
|
|
// CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
|
|
// CHECK2-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]])
|
|
// CHECK2-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8
|
|
// CHECK2-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1
|
|
// CHECK2-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 8
|
|
// CHECK2-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null
|
|
// CHECK2-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]]
|
|
// CHECK2: .select.workers:
|
|
// CHECK2-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1
|
|
// CHECK2-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0
|
|
// CHECK2-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]]
|
|
// CHECK2: .execute.parallel:
|
|
// CHECK2-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]])
|
|
// CHECK2-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)*
|
|
// CHECK2-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]])
|
|
// CHECK2-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]]
|
|
// CHECK2: .terminate.parallel:
|
|
// CHECK2-NEXT: call void @__kmpc_kernel_end_parallel()
|
|
// CHECK2-NEXT: br label [[DOTBARRIER_PARALLEL]]
|
|
// CHECK2: .barrier.parallel:
|
|
// CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
|
|
// CHECK2-NEXT: br label [[DOTAWAIT_WORK]]
|
|
// CHECK2: .exit:
|
|
// CHECK2-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l41
|
|
// CHECK2-SAME: (i64 [[ARGC:%.*]], i32* nonnull align 4 dereferenceable(4) [[B:%.*]], i32* nonnull align 4 dereferenceable(4) [[C:%.*]], i32* [[D:%.*]], i32* nonnull align 4 dereferenceable(4) [[A:%.*]], %class.anon.0* nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK2-NEXT: entry:
|
|
// CHECK2-NEXT: [[ARGC_ADDR:%.*]] = alloca i64, align 8
|
|
// CHECK2-NEXT: [[B_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[D_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[L_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK2-NEXT: [[TMP:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[_TMP1:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[_TMP2:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK2-NEXT: [[L9:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8
|
|
// CHECK2-NEXT: [[_TMP10:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK2-NEXT: [[B11:%.*]] = alloca i32, align 4
|
|
// CHECK2-NEXT: [[_TMP12:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[C13:%.*]] = alloca i32, align 4
|
|
// CHECK2-NEXT: [[_TMP14:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: store i64 [[ARGC]], i64* [[ARGC_ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[B]], i32** [[B_ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[D]], i32** [[D_ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[A]], i32** [[A_ADDR]], align 8
|
|
// CHECK2-NEXT: store %class.anon.0* [[L]], %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK2-NEXT: [[CONV:%.*]] = bitcast i64* [[ARGC_ADDR]] to i32*
|
|
// CHECK2-NEXT: [[TMP0:%.*]] = load i32*, i32** [[B_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP1:%.*]] = load i32*, i32** [[C_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP2:%.*]] = load i32*, i32** [[A_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP3:%.*]] = load %class.anon.0*, %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[TMP0]], i32** [[TMP]], align 8
|
|
// CHECK2-NEXT: store i32* [[TMP1]], i32** [[_TMP1]], align 8
|
|
// CHECK2-NEXT: store %class.anon.0* [[TMP3]], %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK2-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
|
|
// CHECK2-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK2-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
|
|
// CHECK2-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]]
|
|
// CHECK2-NEXT: [[TMP4:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]]
|
|
// CHECK2-NEXT: br i1 [[TMP4]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]]
|
|
// CHECK2: .worker:
|
|
// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l41_worker() #[[ATTR6]]
|
|
// CHECK2-NEXT: br label [[DOTEXIT:%.*]]
|
|
// CHECK2: .mastercheck:
|
|
// CHECK2-NEXT: [[NVPTX_TID3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
|
|
// CHECK2-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK2-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
|
|
// CHECK2-NEXT: [[TMP5:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE5]], 1
|
|
// CHECK2-NEXT: [[TMP6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], 1
|
|
// CHECK2-NEXT: [[TMP7:%.*]] = xor i32 [[TMP5]], -1
|
|
// CHECK2-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP6]], [[TMP7]]
|
|
// CHECK2-NEXT: [[TMP8:%.*]] = icmp eq i32 [[NVPTX_TID3]], [[MASTER_TID]]
|
|
// CHECK2-NEXT: br i1 [[TMP8]], label [[DOTMASTER:%.*]], label [[DOTEXIT]]
|
|
// CHECK2: .master:
|
|
// CHECK2-NEXT: [[NVPTX_NUM_THREADS6:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK2-NEXT: [[NVPTX_WARP_SIZE7:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
|
|
// CHECK2-NEXT: [[THREAD_LIMIT8:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS6]], [[NVPTX_WARP_SIZE7]]
|
|
// CHECK2-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT8]], i16 1)
|
|
// CHECK2-NEXT: call void @__kmpc_data_sharing_init_stack()
|
|
// CHECK2-NEXT: [[TMP9:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK2-NEXT: [[TMP10:%.*]] = bitcast %class.anon.0* [[L9]] to i8*
|
|
// CHECK2-NEXT: [[TMP11:%.*]] = bitcast %class.anon.0* [[TMP9]] to i8*
|
|
// CHECK2-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP10]], i8* align 8 [[TMP11]], i64 40, i1 false)
|
|
// CHECK2-NEXT: store %class.anon.0* [[L9]], %class.anon.0** [[_TMP10]], align 8
|
|
// CHECK2-NEXT: [[TMP12:%.*]] = load i32*, i32** [[TMP]], align 8
|
|
// CHECK2-NEXT: [[TMP13:%.*]] = load i32, i32* [[TMP12]], align 4
|
|
// CHECK2-NEXT: store i32 [[TMP13]], i32* [[B11]], align 4
|
|
// CHECK2-NEXT: store i32* [[B11]], i32** [[_TMP12]], align 8
|
|
// CHECK2-NEXT: [[TMP14:%.*]] = load i32*, i32** [[_TMP1]], align 8
|
|
// CHECK2-NEXT: [[TMP15:%.*]] = load i32, i32* [[TMP14]], align 4
|
|
// CHECK2-NEXT: store i32 [[TMP15]], i32* [[C13]], align 4
|
|
// CHECK2-NEXT: store i32* [[C13]], i32** [[_TMP14]], align 8
|
|
// CHECK2-NEXT: [[TMP16:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP10]], align 8
|
|
// CHECK2-NEXT: [[TMP17:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP16]], i32 0, i32 0
|
|
// CHECK2-NEXT: store i32* [[CONV]], i32** [[TMP17]], align 8
|
|
// CHECK2-NEXT: [[TMP18:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP16]], i32 0, i32 1
|
|
// CHECK2-NEXT: [[TMP19:%.*]] = load i32*, i32** [[_TMP12]], align 8
|
|
// CHECK2-NEXT: store i32* [[TMP19]], i32** [[TMP18]], align 8
|
|
// CHECK2-NEXT: [[TMP20:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP16]], i32 0, i32 2
|
|
// CHECK2-NEXT: [[TMP21:%.*]] = load i32*, i32** [[_TMP14]], align 8
|
|
// CHECK2-NEXT: store i32* [[TMP21]], i32** [[TMP20]], align 8
|
|
// CHECK2-NEXT: [[TMP22:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP16]], i32 0, i32 3
|
|
// CHECK2-NEXT: store i32** [[D_ADDR]], i32*** [[TMP22]], align 8
|
|
// CHECK2-NEXT: [[TMP23:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP16]], i32 0, i32 4
|
|
// CHECK2-NEXT: store i32* [[TMP2]], i32** [[TMP23]], align 8
|
|
// CHECK2-NEXT: [[TMP24:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP10]], align 8
|
|
// CHECK2-NEXT: [[CALL:%.*]] = call i64 @"_ZZ4mainENK3$_0clEv"(%class.anon.0* nonnull dereferenceable(40) [[TMP24]]) #[[ATTR7]]
|
|
// CHECK2-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]]
|
|
// CHECK2: .termination.notifier:
|
|
// CHECK2-NEXT: call void @__kmpc_kernel_deinit(i16 1)
|
|
// CHECK2-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
|
|
// CHECK2-NEXT: br label [[DOTEXIT]]
|
|
// CHECK2: .exit:
|
|
// CHECK2-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l43
|
|
// CHECK2-SAME: (i32* nonnull align 4 dereferenceable(4) [[ARGC:%.*]], i32* nonnull align 4 dereferenceable(4) [[B:%.*]], i32* nonnull align 4 dereferenceable(4) [[C:%.*]], i32* [[D:%.*]], i32* nonnull align 4 dereferenceable(4) [[A:%.*]], %class.anon.0* nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK2-NEXT: entry:
|
|
// CHECK2-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[B_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[D_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[L_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK2-NEXT: [[TMP:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[_TMP1:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[_TMP2:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [6 x i8*], align 8
|
|
// CHECK2-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[B]], i32** [[B_ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[D]], i32** [[D_ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[A]], i32** [[A_ADDR]], align 8
|
|
// CHECK2-NEXT: store %class.anon.0* [[L]], %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP1:%.*]] = load i32*, i32** [[B_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP2:%.*]] = load i32*, i32** [[C_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP3:%.*]] = load i32*, i32** [[A_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP4:%.*]] = load %class.anon.0*, %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[TMP1]], i32** [[TMP]], align 8
|
|
// CHECK2-NEXT: store i32* [[TMP2]], i32** [[_TMP1]], align 8
|
|
// CHECK2-NEXT: store %class.anon.0* [[TMP4]], %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK2-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK2-NEXT: call void @__kmpc_spmd_kernel_init(i32 [[NVPTX_NUM_THREADS]], i16 1)
|
|
// CHECK2-NEXT: call void @__kmpc_data_sharing_init_stack_spmd()
|
|
// CHECK2-NEXT: br label [[DOTEXECUTE:%.*]]
|
|
// CHECK2: .execute:
|
|
// CHECK2-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2]])
|
|
// CHECK2-NEXT: [[TMP6:%.*]] = load i32*, i32** [[TMP]], align 8
|
|
// CHECK2-NEXT: [[TMP7:%.*]] = load i32*, i32** [[_TMP1]], align 8
|
|
// CHECK2-NEXT: [[TMP8:%.*]] = load i32*, i32** [[D_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP9:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK2-NEXT: [[TMP10:%.*]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0
|
|
// CHECK2-NEXT: [[TMP11:%.*]] = bitcast i32* [[TMP0]] to i8*
|
|
// CHECK2-NEXT: store i8* [[TMP11]], i8** [[TMP10]], align 8
|
|
// CHECK2-NEXT: [[TMP12:%.*]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 1
|
|
// CHECK2-NEXT: [[TMP13:%.*]] = bitcast i32* [[TMP6]] to i8*
|
|
// CHECK2-NEXT: store i8* [[TMP13]], i8** [[TMP12]], align 8
|
|
// CHECK2-NEXT: [[TMP14:%.*]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 2
|
|
// CHECK2-NEXT: [[TMP15:%.*]] = bitcast i32* [[TMP7]] to i8*
|
|
// CHECK2-NEXT: store i8* [[TMP15]], i8** [[TMP14]], align 8
|
|
// CHECK2-NEXT: [[TMP16:%.*]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 3
|
|
// CHECK2-NEXT: [[TMP17:%.*]] = bitcast i32* [[TMP8]] to i8*
|
|
// CHECK2-NEXT: store i8* [[TMP17]], i8** [[TMP16]], align 8
|
|
// CHECK2-NEXT: [[TMP18:%.*]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 4
|
|
// CHECK2-NEXT: [[TMP19:%.*]] = bitcast i32* [[TMP3]] to i8*
|
|
// CHECK2-NEXT: store i8* [[TMP19]], i8** [[TMP18]], align 8
|
|
// CHECK2-NEXT: [[TMP20:%.*]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 5
|
|
// CHECK2-NEXT: [[TMP21:%.*]] = bitcast %class.anon.0* [[TMP9]] to i8*
|
|
// CHECK2-NEXT: store i8* [[TMP21]], i8** [[TMP20]], align 8
|
|
// CHECK2-NEXT: [[TMP22:%.*]] = bitcast [6 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
|
|
// CHECK2-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP5]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32*, i32*, i32*, i32*, i32*, %class.anon.0*)* @__omp_outlined__1 to i8*), i8* null, i8** [[TMP22]], i64 6)
|
|
// CHECK2-NEXT: br label [[DOTOMP_DEINIT:%.*]]
|
|
// CHECK2: .omp.deinit:
|
|
// CHECK2-NEXT: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
|
|
// CHECK2-NEXT: br label [[DOTEXIT:%.*]]
|
|
// CHECK2: .exit:
|
|
// CHECK2-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__1
|
|
// CHECK2-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[ARGC:%.*]], i32* nonnull align 4 dereferenceable(4) [[B:%.*]], i32* nonnull align 4 dereferenceable(4) [[C:%.*]], i32* [[D:%.*]], i32* nonnull align 4 dereferenceable(4) [[A:%.*]], %class.anon.0* nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK2-NEXT: entry:
|
|
// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[B_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[D_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[L_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK2-NEXT: [[TMP:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[_TMP1:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[_TMP2:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK2-NEXT: [[L3:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8
|
|
// CHECK2-NEXT: [[_TMP4:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK2-NEXT: [[ARGC5:%.*]] = alloca i32, align 4
|
|
// CHECK2-NEXT: [[B6:%.*]] = alloca i32, align 4
|
|
// CHECK2-NEXT: [[_TMP7:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[C8:%.*]] = alloca i32, align 4
|
|
// CHECK2-NEXT: [[_TMP9:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[A10:%.*]] = alloca i32, align 4
|
|
// CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[B]], i32** [[B_ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[D]], i32** [[D_ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[A]], i32** [[A_ADDR]], align 8
|
|
// CHECK2-NEXT: store %class.anon.0* [[L]], %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP1:%.*]] = load i32*, i32** [[B_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP2:%.*]] = load i32*, i32** [[C_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP3:%.*]] = load i32*, i32** [[A_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP4:%.*]] = load %class.anon.0*, %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[TMP1]], i32** [[TMP]], align 8
|
|
// CHECK2-NEXT: store i32* [[TMP2]], i32** [[_TMP1]], align 8
|
|
// CHECK2-NEXT: store %class.anon.0* [[TMP4]], %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK2-NEXT: [[TMP5:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK2-NEXT: [[TMP6:%.*]] = bitcast %class.anon.0* [[L3]] to i8*
|
|
// CHECK2-NEXT: [[TMP7:%.*]] = bitcast %class.anon.0* [[TMP5]] to i8*
|
|
// CHECK2-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP6]], i8* align 8 [[TMP7]], i64 40, i1 false)
|
|
// CHECK2-NEXT: store %class.anon.0* [[L3]], %class.anon.0** [[_TMP4]], align 8
|
|
// CHECK2-NEXT: [[TMP8:%.*]] = load i32, i32* [[TMP0]], align 4
|
|
// CHECK2-NEXT: store i32 [[TMP8]], i32* [[ARGC5]], align 4
|
|
// CHECK2-NEXT: [[TMP9:%.*]] = load i32*, i32** [[TMP]], align 8
|
|
// CHECK2-NEXT: [[TMP10:%.*]] = load i32, i32* [[TMP9]], align 4
|
|
// CHECK2-NEXT: store i32 [[TMP10]], i32* [[B6]], align 4
|
|
// CHECK2-NEXT: store i32* [[B6]], i32** [[_TMP7]], align 8
|
|
// CHECK2-NEXT: [[TMP11:%.*]] = load i32*, i32** [[_TMP1]], align 8
|
|
// CHECK2-NEXT: [[TMP12:%.*]] = load i32, i32* [[TMP11]], align 4
|
|
// CHECK2-NEXT: store i32 [[TMP12]], i32* [[C8]], align 4
|
|
// CHECK2-NEXT: store i32* [[C8]], i32** [[_TMP9]], align 8
|
|
// CHECK2-NEXT: [[TMP13:%.*]] = load i32, i32* [[TMP3]], align 4
|
|
// CHECK2-NEXT: store i32 [[TMP13]], i32* [[A10]], align 4
|
|
// CHECK2-NEXT: [[TMP14:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP4]], align 8
|
|
// CHECK2-NEXT: [[TMP15:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP14]], i32 0, i32 0
|
|
// CHECK2-NEXT: store i32* [[ARGC5]], i32** [[TMP15]], align 8
|
|
// CHECK2-NEXT: [[TMP16:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP14]], i32 0, i32 1
|
|
// CHECK2-NEXT: [[TMP17:%.*]] = load i32*, i32** [[_TMP7]], align 8
|
|
// CHECK2-NEXT: store i32* [[TMP17]], i32** [[TMP16]], align 8
|
|
// CHECK2-NEXT: [[TMP18:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP14]], i32 0, i32 2
|
|
// CHECK2-NEXT: [[TMP19:%.*]] = load i32*, i32** [[_TMP9]], align 8
|
|
// CHECK2-NEXT: store i32* [[TMP19]], i32** [[TMP18]], align 8
|
|
// CHECK2-NEXT: [[TMP20:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP14]], i32 0, i32 3
|
|
// CHECK2-NEXT: store i32** [[D_ADDR]], i32*** [[TMP20]], align 8
|
|
// CHECK2-NEXT: [[TMP21:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP14]], i32 0, i32 4
|
|
// CHECK2-NEXT: store i32* [[A10]], i32** [[TMP21]], align 8
|
|
// CHECK2-NEXT: [[TMP22:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP4]], align 8
|
|
// CHECK2-NEXT: [[CALL:%.*]] = call i64 @"_ZZ4mainENK3$_0clEv"(%class.anon.0* nonnull dereferenceable(40) [[TMP22]]) #[[ATTR7]]
|
|
// CHECK2-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l18
|
|
// CHECK2-SAME: (%class.anon* nonnull align 8 dereferenceable(8) [[T:%.*]]) #[[ATTR1]] {
|
|
// CHECK2-NEXT: entry:
|
|
// CHECK2-NEXT: [[T_ADDR:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK2-NEXT: [[TMP:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8
|
|
// CHECK2-NEXT: store %class.anon* [[T]], %class.anon** [[T_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP0:%.*]] = load %class.anon*, %class.anon** [[T_ADDR]], align 8
|
|
// CHECK2-NEXT: store %class.anon* [[TMP0]], %class.anon** [[TMP]], align 8
|
|
// CHECK2-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK2-NEXT: call void @__kmpc_spmd_kernel_init(i32 [[NVPTX_NUM_THREADS]], i16 1)
|
|
// CHECK2-NEXT: call void @__kmpc_data_sharing_init_stack_spmd()
|
|
// CHECK2-NEXT: br label [[DOTEXECUTE:%.*]]
|
|
// CHECK2: .execute:
|
|
// CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2]])
|
|
// CHECK2-NEXT: [[TMP2:%.*]] = load %class.anon*, %class.anon** [[TMP]], align 8
|
|
// CHECK2-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0
|
|
// CHECK2-NEXT: [[TMP4:%.*]] = bitcast %class.anon* [[TMP2]] to i8*
|
|
// CHECK2-NEXT: store i8* [[TMP4]], i8** [[TMP3]], align 8
|
|
// CHECK2-NEXT: [[TMP5:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
|
|
// CHECK2-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, %class.anon*)* @__omp_outlined__2 to i8*), i8* null, i8** [[TMP5]], i64 1)
|
|
// CHECK2-NEXT: br label [[DOTOMP_DEINIT:%.*]]
|
|
// CHECK2: .omp.deinit:
|
|
// CHECK2-NEXT: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
|
|
// CHECK2-NEXT: br label [[DOTEXIT:%.*]]
|
|
// CHECK2: .exit:
|
|
// CHECK2-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__2
|
|
// CHECK2-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], %class.anon* nonnull align 8 dereferenceable(8) [[T:%.*]]) #[[ATTR1]] {
|
|
// CHECK2-NEXT: entry:
|
|
// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK2-NEXT: [[T_ADDR:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK2-NEXT: [[TMP:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK2-NEXT: [[T1:%.*]] = alloca [[CLASS_ANON:%.*]], align 8
|
|
// CHECK2-NEXT: [[_TMP2:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
|
|
// CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
|
|
// CHECK2-NEXT: store %class.anon* [[T]], %class.anon** [[T_ADDR]], align 8
|
|
// CHECK2-NEXT: [[TMP0:%.*]] = load %class.anon*, %class.anon** [[T_ADDR]], align 8
|
|
// CHECK2-NEXT: store %class.anon* [[TMP0]], %class.anon** [[TMP]], align 8
|
|
// CHECK2-NEXT: [[TMP1:%.*]] = load %class.anon*, %class.anon** [[TMP]], align 8
|
|
// CHECK2-NEXT: [[TMP2:%.*]] = bitcast %class.anon* [[T1]] to i8*
|
|
// CHECK2-NEXT: [[TMP3:%.*]] = bitcast %class.anon* [[TMP1]] to i8*
|
|
// CHECK2-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP2]], i8* align 8 [[TMP3]], i64 8, i1 false)
|
|
// CHECK2-NEXT: store %class.anon* [[T1]], %class.anon** [[_TMP2]], align 8
|
|
// CHECK2-NEXT: [[TMP4:%.*]] = load %class.anon*, %class.anon** [[_TMP2]], align 8
|
|
// CHECK2-NEXT: [[TMP5:%.*]] = load %class.anon*, %class.anon** [[_TMP2]], align 8
|
|
// CHECK2-NEXT: [[CALL:%.*]] = call i32 @_ZZN1S3fooEvENKUlvE_clEv(%class.anon* nonnull dereferenceable(8) [[TMP5]]) #[[ATTR7]]
|
|
// CHECK2-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l41_worker
|
|
// CHECK3-SAME: () #[[ATTR0:[0-9]+]] {
|
|
// CHECK3-NEXT: entry:
|
|
// CHECK3-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 8
|
|
// CHECK3-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1
|
|
// CHECK3-NEXT: store i8* null, i8** [[WORK_FN]], align 8
|
|
// CHECK3-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1
|
|
// CHECK3-NEXT: br label [[DOTAWAIT_WORK:%.*]]
|
|
// CHECK3: .await.work:
|
|
// CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
|
|
// CHECK3-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]])
|
|
// CHECK3-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8
|
|
// CHECK3-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1
|
|
// CHECK3-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 8
|
|
// CHECK3-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null
|
|
// CHECK3-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]]
|
|
// CHECK3: .select.workers:
|
|
// CHECK3-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1
|
|
// CHECK3-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0
|
|
// CHECK3-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]]
|
|
// CHECK3: .execute.parallel:
|
|
// CHECK3-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]])
|
|
// CHECK3-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)*
|
|
// CHECK3-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]])
|
|
// CHECK3-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]]
|
|
// CHECK3: .terminate.parallel:
|
|
// CHECK3-NEXT: call void @__kmpc_kernel_end_parallel()
|
|
// CHECK3-NEXT: br label [[DOTBARRIER_PARALLEL]]
|
|
// CHECK3: .barrier.parallel:
|
|
// CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
|
|
// CHECK3-NEXT: br label [[DOTAWAIT_WORK]]
|
|
// CHECK3: .exit:
|
|
// CHECK3-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l41
|
|
// CHECK3-SAME: (i64 [[ARGC:%.*]], i32* nonnull align 4 dereferenceable(4) [[B:%.*]], i32* nonnull align 4 dereferenceable(4) [[C:%.*]], i32* [[D:%.*]], i32* nonnull align 4 dereferenceable(4) [[A:%.*]], %class.anon* nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR1:[0-9]+]] {
|
|
// CHECK3-NEXT: entry:
|
|
// CHECK3-NEXT: [[ARGC_ADDR:%.*]] = alloca i64, align 8
|
|
// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[D_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[L_ADDR:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK3-NEXT: [[TMP:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[_TMP1:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[_TMP2:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK3-NEXT: [[L9:%.*]] = alloca [[CLASS_ANON:%.*]], align 8
|
|
// CHECK3-NEXT: [[_TMP10:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK3-NEXT: [[B11:%.*]] = alloca i32, align 4
|
|
// CHECK3-NEXT: [[_TMP12:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[C13:%.*]] = alloca i32, align 4
|
|
// CHECK3-NEXT: [[_TMP14:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: store i64 [[ARGC]], i64* [[ARGC_ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[B]], i32** [[B_ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[D]], i32** [[D_ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[A]], i32** [[A_ADDR]], align 8
|
|
// CHECK3-NEXT: store %class.anon* [[L]], %class.anon** [[L_ADDR]], align 8
|
|
// CHECK3-NEXT: [[CONV:%.*]] = bitcast i64* [[ARGC_ADDR]] to i32*
|
|
// CHECK3-NEXT: [[TMP0:%.*]] = load i32*, i32** [[B_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP1:%.*]] = load i32*, i32** [[C_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP2:%.*]] = load i32*, i32** [[A_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP3:%.*]] = load %class.anon*, %class.anon** [[L_ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[TMP0]], i32** [[TMP]], align 8
|
|
// CHECK3-NEXT: store i32* [[TMP1]], i32** [[_TMP1]], align 8
|
|
// CHECK3-NEXT: store %class.anon* [[TMP3]], %class.anon** [[_TMP2]], align 8
|
|
// CHECK3-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
|
|
// CHECK3-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK3-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
|
|
// CHECK3-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]]
|
|
// CHECK3-NEXT: [[TMP4:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]]
|
|
// CHECK3-NEXT: br i1 [[TMP4]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]]
|
|
// CHECK3: .worker:
|
|
// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l41_worker() #[[ATTR6:[0-9]+]]
|
|
// CHECK3-NEXT: br label [[DOTEXIT:%.*]]
|
|
// CHECK3: .mastercheck:
|
|
// CHECK3-NEXT: [[NVPTX_TID3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
|
|
// CHECK3-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK3-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
|
|
// CHECK3-NEXT: [[TMP5:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE5]], 1
|
|
// CHECK3-NEXT: [[TMP6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], 1
|
|
// CHECK3-NEXT: [[TMP7:%.*]] = xor i32 [[TMP5]], -1
|
|
// CHECK3-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP6]], [[TMP7]]
|
|
// CHECK3-NEXT: [[TMP8:%.*]] = icmp eq i32 [[NVPTX_TID3]], [[MASTER_TID]]
|
|
// CHECK3-NEXT: br i1 [[TMP8]], label [[DOTMASTER:%.*]], label [[DOTEXIT]]
|
|
// CHECK3: .master:
|
|
// CHECK3-NEXT: [[NVPTX_NUM_THREADS6:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK3-NEXT: [[NVPTX_WARP_SIZE7:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
|
|
// CHECK3-NEXT: [[THREAD_LIMIT8:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS6]], [[NVPTX_WARP_SIZE7]]
|
|
// CHECK3-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT8]], i16 1)
|
|
// CHECK3-NEXT: call void @__kmpc_data_sharing_init_stack()
|
|
// CHECK3-NEXT: [[TMP9:%.*]] = load %class.anon*, %class.anon** [[_TMP2]], align 8
|
|
// CHECK3-NEXT: [[TMP10:%.*]] = bitcast %class.anon* [[L9]] to i8*
|
|
// CHECK3-NEXT: [[TMP11:%.*]] = bitcast %class.anon* [[TMP9]] to i8*
|
|
// CHECK3-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP10]], i8* align 8 [[TMP11]], i64 40, i1 false)
|
|
// CHECK3-NEXT: store %class.anon* [[L9]], %class.anon** [[_TMP10]], align 8
|
|
// CHECK3-NEXT: [[TMP12:%.*]] = load i32*, i32** [[TMP]], align 8
|
|
// CHECK3-NEXT: [[TMP13:%.*]] = load i32, i32* [[TMP12]], align 4
|
|
// CHECK3-NEXT: store i32 [[TMP13]], i32* [[B11]], align 4
|
|
// CHECK3-NEXT: store i32* [[B11]], i32** [[_TMP12]], align 8
|
|
// CHECK3-NEXT: [[TMP14:%.*]] = load i32*, i32** [[_TMP1]], align 8
|
|
// CHECK3-NEXT: [[TMP15:%.*]] = load i32, i32* [[TMP14]], align 4
|
|
// CHECK3-NEXT: store i32 [[TMP15]], i32* [[C13]], align 4
|
|
// CHECK3-NEXT: store i32* [[C13]], i32** [[_TMP14]], align 8
|
|
// CHECK3-NEXT: [[TMP16:%.*]] = load %class.anon*, %class.anon** [[_TMP10]], align 8
|
|
// CHECK3-NEXT: [[TMP17:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP16]], i32 0, i32 0
|
|
// CHECK3-NEXT: store i32* [[CONV]], i32** [[TMP17]], align 8
|
|
// CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP16]], i32 0, i32 1
|
|
// CHECK3-NEXT: [[TMP19:%.*]] = load i32*, i32** [[_TMP12]], align 8
|
|
// CHECK3-NEXT: store i32* [[TMP19]], i32** [[TMP18]], align 8
|
|
// CHECK3-NEXT: [[TMP20:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP16]], i32 0, i32 2
|
|
// CHECK3-NEXT: [[TMP21:%.*]] = load i32*, i32** [[_TMP14]], align 8
|
|
// CHECK3-NEXT: store i32* [[TMP21]], i32** [[TMP20]], align 8
|
|
// CHECK3-NEXT: [[TMP22:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP16]], i32 0, i32 3
|
|
// CHECK3-NEXT: store i32** [[D_ADDR]], i32*** [[TMP22]], align 8
|
|
// CHECK3-NEXT: [[TMP23:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP16]], i32 0, i32 4
|
|
// CHECK3-NEXT: store i32* [[TMP2]], i32** [[TMP23]], align 8
|
|
// CHECK3-NEXT: [[TMP24:%.*]] = load %class.anon*, %class.anon** [[_TMP10]], align 8
|
|
// CHECK3-NEXT: [[CALL:%.*]] = call i64 @"_ZZ4mainENK3$_0clEv"(%class.anon* nonnull dereferenceable(40) [[TMP24]]) #[[ATTR7:[0-9]+]]
|
|
// CHECK3-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]]
|
|
// CHECK3: .termination.notifier:
|
|
// CHECK3-NEXT: call void @__kmpc_kernel_deinit(i16 1)
|
|
// CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
|
|
// CHECK3-NEXT: br label [[DOTEXIT]]
|
|
// CHECK3: .exit:
|
|
// CHECK3-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l43
|
|
// CHECK3-SAME: (i32* nonnull align 4 dereferenceable(4) [[ARGC:%.*]], i32* nonnull align 4 dereferenceable(4) [[B:%.*]], i32* nonnull align 4 dereferenceable(4) [[C:%.*]], i32* [[D:%.*]], i32* nonnull align 4 dereferenceable(4) [[A:%.*]], %class.anon* nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK3-NEXT: entry:
|
|
// CHECK3-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[D_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[L_ADDR:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK3-NEXT: [[TMP:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[_TMP1:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[_TMP2:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK3-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [6 x i8*], align 8
|
|
// CHECK3-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[B]], i32** [[B_ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[D]], i32** [[D_ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[A]], i32** [[A_ADDR]], align 8
|
|
// CHECK3-NEXT: store %class.anon* [[L]], %class.anon** [[L_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP1:%.*]] = load i32*, i32** [[B_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP2:%.*]] = load i32*, i32** [[C_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP3:%.*]] = load i32*, i32** [[A_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP4:%.*]] = load %class.anon*, %class.anon** [[L_ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[TMP1]], i32** [[TMP]], align 8
|
|
// CHECK3-NEXT: store i32* [[TMP2]], i32** [[_TMP1]], align 8
|
|
// CHECK3-NEXT: store %class.anon* [[TMP4]], %class.anon** [[_TMP2]], align 8
|
|
// CHECK3-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK3-NEXT: call void @__kmpc_spmd_kernel_init(i32 [[NVPTX_NUM_THREADS]], i16 1)
|
|
// CHECK3-NEXT: call void @__kmpc_data_sharing_init_stack_spmd()
|
|
// CHECK3-NEXT: br label [[DOTEXECUTE:%.*]]
|
|
// CHECK3: .execute:
|
|
// CHECK3-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2:[0-9]+]])
|
|
// CHECK3-NEXT: [[TMP6:%.*]] = load i32*, i32** [[TMP]], align 8
|
|
// CHECK3-NEXT: [[TMP7:%.*]] = load i32*, i32** [[_TMP1]], align 8
|
|
// CHECK3-NEXT: [[TMP8:%.*]] = load i32*, i32** [[D_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP9:%.*]] = load %class.anon*, %class.anon** [[_TMP2]], align 8
|
|
// CHECK3-NEXT: [[TMP10:%.*]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0
|
|
// CHECK3-NEXT: [[TMP11:%.*]] = bitcast i32* [[TMP0]] to i8*
|
|
// CHECK3-NEXT: store i8* [[TMP11]], i8** [[TMP10]], align 8
|
|
// CHECK3-NEXT: [[TMP12:%.*]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 1
|
|
// CHECK3-NEXT: [[TMP13:%.*]] = bitcast i32* [[TMP6]] to i8*
|
|
// CHECK3-NEXT: store i8* [[TMP13]], i8** [[TMP12]], align 8
|
|
// CHECK3-NEXT: [[TMP14:%.*]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 2
|
|
// CHECK3-NEXT: [[TMP15:%.*]] = bitcast i32* [[TMP7]] to i8*
|
|
// CHECK3-NEXT: store i8* [[TMP15]], i8** [[TMP14]], align 8
|
|
// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 3
|
|
// CHECK3-NEXT: [[TMP17:%.*]] = bitcast i32* [[TMP8]] to i8*
|
|
// CHECK3-NEXT: store i8* [[TMP17]], i8** [[TMP16]], align 8
|
|
// CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 4
|
|
// CHECK3-NEXT: [[TMP19:%.*]] = bitcast i32* [[TMP3]] to i8*
|
|
// CHECK3-NEXT: store i8* [[TMP19]], i8** [[TMP18]], align 8
|
|
// CHECK3-NEXT: [[TMP20:%.*]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 5
|
|
// CHECK3-NEXT: [[TMP21:%.*]] = bitcast %class.anon* [[TMP9]] to i8*
|
|
// CHECK3-NEXT: store i8* [[TMP21]], i8** [[TMP20]], align 8
|
|
// CHECK3-NEXT: [[TMP22:%.*]] = bitcast [6 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
|
|
// CHECK3-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP5]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32*, i32*, i32*, i32*, i32*, %class.anon*)* @__omp_outlined__ to i8*), i8* null, i8** [[TMP22]], i64 6)
|
|
// CHECK3-NEXT: br label [[DOTOMP_DEINIT:%.*]]
|
|
// CHECK3: .omp.deinit:
|
|
// CHECK3-NEXT: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
|
|
// CHECK3-NEXT: br label [[DOTEXIT:%.*]]
|
|
// CHECK3: .exit:
|
|
// CHECK3-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__
|
|
// CHECK3-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[ARGC:%.*]], i32* nonnull align 4 dereferenceable(4) [[B:%.*]], i32* nonnull align 4 dereferenceable(4) [[C:%.*]], i32* [[D:%.*]], i32* nonnull align 4 dereferenceable(4) [[A:%.*]], %class.anon* nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK3-NEXT: entry:
|
|
// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[B_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[D_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[L_ADDR:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK3-NEXT: [[TMP:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[_TMP1:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[_TMP2:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK3-NEXT: [[L3:%.*]] = alloca [[CLASS_ANON:%.*]], align 8
|
|
// CHECK3-NEXT: [[_TMP4:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK3-NEXT: [[ARGC5:%.*]] = alloca i32, align 4
|
|
// CHECK3-NEXT: [[B6:%.*]] = alloca i32, align 4
|
|
// CHECK3-NEXT: [[_TMP7:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[C8:%.*]] = alloca i32, align 4
|
|
// CHECK3-NEXT: [[_TMP9:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[A10:%.*]] = alloca i32, align 4
|
|
// CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[B]], i32** [[B_ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[D]], i32** [[D_ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[A]], i32** [[A_ADDR]], align 8
|
|
// CHECK3-NEXT: store %class.anon* [[L]], %class.anon** [[L_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP1:%.*]] = load i32*, i32** [[B_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP2:%.*]] = load i32*, i32** [[C_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP3:%.*]] = load i32*, i32** [[A_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP4:%.*]] = load %class.anon*, %class.anon** [[L_ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[TMP1]], i32** [[TMP]], align 8
|
|
// CHECK3-NEXT: store i32* [[TMP2]], i32** [[_TMP1]], align 8
|
|
// CHECK3-NEXT: store %class.anon* [[TMP4]], %class.anon** [[_TMP2]], align 8
|
|
// CHECK3-NEXT: [[TMP5:%.*]] = load %class.anon*, %class.anon** [[_TMP2]], align 8
|
|
// CHECK3-NEXT: [[TMP6:%.*]] = bitcast %class.anon* [[L3]] to i8*
|
|
// CHECK3-NEXT: [[TMP7:%.*]] = bitcast %class.anon* [[TMP5]] to i8*
|
|
// CHECK3-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP6]], i8* align 8 [[TMP7]], i64 40, i1 false)
|
|
// CHECK3-NEXT: store %class.anon* [[L3]], %class.anon** [[_TMP4]], align 8
|
|
// CHECK3-NEXT: [[TMP8:%.*]] = load i32, i32* [[TMP0]], align 4
|
|
// CHECK3-NEXT: store i32 [[TMP8]], i32* [[ARGC5]], align 4
|
|
// CHECK3-NEXT: [[TMP9:%.*]] = load i32*, i32** [[TMP]], align 8
|
|
// CHECK3-NEXT: [[TMP10:%.*]] = load i32, i32* [[TMP9]], align 4
|
|
// CHECK3-NEXT: store i32 [[TMP10]], i32* [[B6]], align 4
|
|
// CHECK3-NEXT: store i32* [[B6]], i32** [[_TMP7]], align 8
|
|
// CHECK3-NEXT: [[TMP11:%.*]] = load i32*, i32** [[_TMP1]], align 8
|
|
// CHECK3-NEXT: [[TMP12:%.*]] = load i32, i32* [[TMP11]], align 4
|
|
// CHECK3-NEXT: store i32 [[TMP12]], i32* [[C8]], align 4
|
|
// CHECK3-NEXT: store i32* [[C8]], i32** [[_TMP9]], align 8
|
|
// CHECK3-NEXT: [[TMP13:%.*]] = load i32, i32* [[TMP3]], align 4
|
|
// CHECK3-NEXT: store i32 [[TMP13]], i32* [[A10]], align 4
|
|
// CHECK3-NEXT: [[TMP14:%.*]] = load %class.anon*, %class.anon** [[_TMP4]], align 8
|
|
// CHECK3-NEXT: [[TMP15:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP14]], i32 0, i32 0
|
|
// CHECK3-NEXT: store i32* [[ARGC5]], i32** [[TMP15]], align 8
|
|
// CHECK3-NEXT: [[TMP16:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP14]], i32 0, i32 1
|
|
// CHECK3-NEXT: [[TMP17:%.*]] = load i32*, i32** [[_TMP7]], align 8
|
|
// CHECK3-NEXT: store i32* [[TMP17]], i32** [[TMP16]], align 8
|
|
// CHECK3-NEXT: [[TMP18:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP14]], i32 0, i32 2
|
|
// CHECK3-NEXT: [[TMP19:%.*]] = load i32*, i32** [[_TMP9]], align 8
|
|
// CHECK3-NEXT: store i32* [[TMP19]], i32** [[TMP18]], align 8
|
|
// CHECK3-NEXT: [[TMP20:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP14]], i32 0, i32 3
|
|
// CHECK3-NEXT: store i32** [[D_ADDR]], i32*** [[TMP20]], align 8
|
|
// CHECK3-NEXT: [[TMP21:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP14]], i32 0, i32 4
|
|
// CHECK3-NEXT: store i32* [[A10]], i32** [[TMP21]], align 8
|
|
// CHECK3-NEXT: [[TMP22:%.*]] = load %class.anon*, %class.anon** [[_TMP4]], align 8
|
|
// CHECK3-NEXT: [[CALL:%.*]] = call i64 @"_ZZ4mainENK3$_0clEv"(%class.anon* nonnull dereferenceable(40) [[TMP22]]) #[[ATTR7]]
|
|
// CHECK3-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l27_worker
|
|
// CHECK3-SAME: () #[[ATTR0]] {
|
|
// CHECK3-NEXT: entry:
|
|
// CHECK3-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 8
|
|
// CHECK3-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1
|
|
// CHECK3-NEXT: store i8* null, i8** [[WORK_FN]], align 8
|
|
// CHECK3-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1
|
|
// CHECK3-NEXT: br label [[DOTAWAIT_WORK:%.*]]
|
|
// CHECK3: .await.work:
|
|
// CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
|
|
// CHECK3-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]])
|
|
// CHECK3-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8
|
|
// CHECK3-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1
|
|
// CHECK3-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 8
|
|
// CHECK3-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null
|
|
// CHECK3-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]]
|
|
// CHECK3: .select.workers:
|
|
// CHECK3-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1
|
|
// CHECK3-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0
|
|
// CHECK3-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]]
|
|
// CHECK3: .execute.parallel:
|
|
// CHECK3-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]])
|
|
// CHECK3-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)*
|
|
// CHECK3-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]])
|
|
// CHECK3-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]]
|
|
// CHECK3: .terminate.parallel:
|
|
// CHECK3-NEXT: call void @__kmpc_kernel_end_parallel()
|
|
// CHECK3-NEXT: br label [[DOTBARRIER_PARALLEL]]
|
|
// CHECK3: .barrier.parallel:
|
|
// CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
|
|
// CHECK3-NEXT: br label [[DOTAWAIT_WORK]]
|
|
// CHECK3: .exit:
|
|
// CHECK3-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l27
|
|
// CHECK3-SAME: (%struct.S* [[THIS:%.*]], %class.anon.0* nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK3-NEXT: entry:
|
|
// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.S*, align 8
|
|
// CHECK3-NEXT: [[L_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK3-NEXT: [[TMP:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK3-NEXT: [[L7:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8
|
|
// CHECK3-NEXT: [[_TMP8:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK3-NEXT: store %struct.S* [[THIS]], %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK3-NEXT: store %class.anon.0* [[L]], %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP0:%.*]] = load %struct.S*, %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP1:%.*]] = load %class.anon.0*, %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK3-NEXT: store %class.anon.0* [[TMP1]], %class.anon.0** [[TMP]], align 8
|
|
// CHECK3-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
|
|
// CHECK3-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK3-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
|
|
// CHECK3-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]]
|
|
// CHECK3-NEXT: [[TMP2:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]]
|
|
// CHECK3-NEXT: br i1 [[TMP2]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]]
|
|
// CHECK3: .worker:
|
|
// CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l27_worker() #[[ATTR6]]
|
|
// CHECK3-NEXT: br label [[DOTEXIT:%.*]]
|
|
// CHECK3: .mastercheck:
|
|
// CHECK3-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
|
|
// CHECK3-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK3-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
|
|
// CHECK3-NEXT: [[TMP3:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1
|
|
// CHECK3-NEXT: [[TMP4:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1
|
|
// CHECK3-NEXT: [[TMP5:%.*]] = xor i32 [[TMP3]], -1
|
|
// CHECK3-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP4]], [[TMP5]]
|
|
// CHECK3-NEXT: [[TMP6:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]]
|
|
// CHECK3-NEXT: br i1 [[TMP6]], label [[DOTMASTER:%.*]], label [[DOTEXIT]]
|
|
// CHECK3: .master:
|
|
// CHECK3-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK3-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
|
|
// CHECK3-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]]
|
|
// CHECK3-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1)
|
|
// CHECK3-NEXT: call void @__kmpc_data_sharing_init_stack()
|
|
// CHECK3-NEXT: [[TMP7:%.*]] = load %class.anon.0*, %class.anon.0** [[TMP]], align 8
|
|
// CHECK3-NEXT: [[TMP8:%.*]] = bitcast %class.anon.0* [[L7]] to i8*
|
|
// CHECK3-NEXT: [[TMP9:%.*]] = bitcast %class.anon.0* [[TMP7]] to i8*
|
|
// CHECK3-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP8]], i8* align 8 [[TMP9]], i64 8, i1 false)
|
|
// CHECK3-NEXT: store %class.anon.0* [[L7]], %class.anon.0** [[_TMP8]], align 8
|
|
// CHECK3-NEXT: [[TMP10:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP8]], align 8
|
|
// CHECK3-NEXT: [[TMP11:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP10]], i32 0, i32 0
|
|
// CHECK3-NEXT: store %struct.S* [[TMP0]], %struct.S** [[TMP11]], align 8
|
|
// CHECK3-NEXT: [[TMP12:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP8]], align 8
|
|
// CHECK3-NEXT: [[CALL:%.*]] = call i32 @_ZZN1S3fooEvENKUlvE_clEv(%class.anon.0* nonnull dereferenceable(8) [[TMP12]]) #[[ATTR7]]
|
|
// CHECK3-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]]
|
|
// CHECK3: .termination.notifier:
|
|
// CHECK3-NEXT: call void @__kmpc_kernel_deinit(i16 1)
|
|
// CHECK3-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
|
|
// CHECK3-NEXT: br label [[DOTEXIT]]
|
|
// CHECK3: .exit:
|
|
// CHECK3-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK3-LABEL: define {{[^@]+}}@_ZZN1S3fooEvENKUlvE_clEv
|
|
// CHECK3-SAME: (%class.anon.0* nonnull dereferenceable(8) [[THIS:%.*]]) #[[ATTR4:[0-9]+]] comdat align 2 {
|
|
// CHECK3-NEXT: entry:
|
|
// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK3-NEXT: store %class.anon.0* [[THIS]], %class.anon.0** [[THIS_ADDR]], align 8
|
|
// CHECK3-NEXT: [[THIS1:%.*]] = load %class.anon.0*, %class.anon.0** [[THIS_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP0:%.*]] = getelementptr inbounds [[CLASS_ANON_0:%.*]], %class.anon.0* [[THIS1]], i32 0, i32 0
|
|
// CHECK3-NEXT: [[TMP1:%.*]] = load %struct.S*, %struct.S** [[TMP0]], align 8
|
|
// CHECK3-NEXT: [[A:%.*]] = getelementptr inbounds [[STRUCT_S:%.*]], %struct.S* [[TMP1]], i32 0, i32 0
|
|
// CHECK3-NEXT: [[TMP2:%.*]] = load i32, i32* [[A]], align 4
|
|
// CHECK3-NEXT: ret i32 [[TMP2]]
|
|
//
|
|
//
|
|
// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l29
|
|
// CHECK3-SAME: (%struct.S* [[THIS:%.*]], %class.anon.0* nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK3-NEXT: entry:
|
|
// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.S*, align 8
|
|
// CHECK3-NEXT: [[L_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK3-NEXT: [[TMP:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK3-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [2 x i8*], align 8
|
|
// CHECK3-NEXT: store %struct.S* [[THIS]], %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK3-NEXT: store %class.anon.0* [[L]], %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP0:%.*]] = load %struct.S*, %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP1:%.*]] = load %class.anon.0*, %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK3-NEXT: store %class.anon.0* [[TMP1]], %class.anon.0** [[TMP]], align 8
|
|
// CHECK3-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK3-NEXT: call void @__kmpc_spmd_kernel_init(i32 [[NVPTX_NUM_THREADS]], i16 1)
|
|
// CHECK3-NEXT: call void @__kmpc_data_sharing_init_stack_spmd()
|
|
// CHECK3-NEXT: br label [[DOTEXECUTE:%.*]]
|
|
// CHECK3: .execute:
|
|
// CHECK3-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2]])
|
|
// CHECK3-NEXT: [[TMP3:%.*]] = load %class.anon.0*, %class.anon.0** [[TMP]], align 8
|
|
// CHECK3-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0
|
|
// CHECK3-NEXT: [[TMP5:%.*]] = bitcast %struct.S* [[TMP0]] to i8*
|
|
// CHECK3-NEXT: store i8* [[TMP5]], i8** [[TMP4]], align 8
|
|
// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 1
|
|
// CHECK3-NEXT: [[TMP7:%.*]] = bitcast %class.anon.0* [[TMP3]] to i8*
|
|
// CHECK3-NEXT: store i8* [[TMP7]], i8** [[TMP6]], align 8
|
|
// CHECK3-NEXT: [[TMP8:%.*]] = bitcast [2 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
|
|
// CHECK3-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP2]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, %struct.S*, %class.anon.0*)* @__omp_outlined__1 to i8*), i8* null, i8** [[TMP8]], i64 2)
|
|
// CHECK3-NEXT: br label [[DOTOMP_DEINIT:%.*]]
|
|
// CHECK3: .omp.deinit:
|
|
// CHECK3-NEXT: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
|
|
// CHECK3-NEXT: br label [[DOTEXIT:%.*]]
|
|
// CHECK3: .exit:
|
|
// CHECK3-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__1
|
|
// CHECK3-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], %struct.S* [[THIS:%.*]], %class.anon.0* nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK3-NEXT: entry:
|
|
// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.S*, align 8
|
|
// CHECK3-NEXT: [[L_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK3-NEXT: [[TMP:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK3-NEXT: [[L1:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8
|
|
// CHECK3-NEXT: [[_TMP2:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
|
|
// CHECK3-NEXT: store %struct.S* [[THIS]], %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK3-NEXT: store %class.anon.0* [[L]], %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP0:%.*]] = load %struct.S*, %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP1:%.*]] = load %class.anon.0*, %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK3-NEXT: store %class.anon.0* [[TMP1]], %class.anon.0** [[TMP]], align 8
|
|
// CHECK3-NEXT: [[TMP2:%.*]] = load %class.anon.0*, %class.anon.0** [[TMP]], align 8
|
|
// CHECK3-NEXT: [[TMP3:%.*]] = bitcast %class.anon.0* [[L1]] to i8*
|
|
// CHECK3-NEXT: [[TMP4:%.*]] = bitcast %class.anon.0* [[TMP2]] to i8*
|
|
// CHECK3-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP3]], i8* align 8 [[TMP4]], i64 8, i1 false)
|
|
// CHECK3-NEXT: store %class.anon.0* [[L1]], %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK3-NEXT: [[TMP5:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK3-NEXT: [[TMP6:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP5]], i32 0, i32 0
|
|
// CHECK3-NEXT: store %struct.S* [[TMP0]], %struct.S** [[TMP6]], align 8
|
|
// CHECK3-NEXT: [[TMP7:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK3-NEXT: [[CALL:%.*]] = call i32 @_ZZN1S3fooEvENKUlvE_clEv(%class.anon.0* nonnull dereferenceable(8) [[TMP7]]) #[[ATTR7]]
|
|
// CHECK3-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l18
|
|
// CHECK3-SAME: (%class.anon.0* nonnull align 8 dereferenceable(8) [[T:%.*]]) #[[ATTR1]] {
|
|
// CHECK3-NEXT: entry:
|
|
// CHECK3-NEXT: [[T_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK3-NEXT: [[TMP:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK3-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8
|
|
// CHECK3-NEXT: store %class.anon.0* [[T]], %class.anon.0** [[T_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP0:%.*]] = load %class.anon.0*, %class.anon.0** [[T_ADDR]], align 8
|
|
// CHECK3-NEXT: store %class.anon.0* [[TMP0]], %class.anon.0** [[TMP]], align 8
|
|
// CHECK3-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK3-NEXT: call void @__kmpc_spmd_kernel_init(i32 [[NVPTX_NUM_THREADS]], i16 1)
|
|
// CHECK3-NEXT: call void @__kmpc_data_sharing_init_stack_spmd()
|
|
// CHECK3-NEXT: br label [[DOTEXECUTE:%.*]]
|
|
// CHECK3: .execute:
|
|
// CHECK3-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2]])
|
|
// CHECK3-NEXT: [[TMP2:%.*]] = load %class.anon.0*, %class.anon.0** [[TMP]], align 8
|
|
// CHECK3-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0
|
|
// CHECK3-NEXT: [[TMP4:%.*]] = bitcast %class.anon.0* [[TMP2]] to i8*
|
|
// CHECK3-NEXT: store i8* [[TMP4]], i8** [[TMP3]], align 8
|
|
// CHECK3-NEXT: [[TMP5:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
|
|
// CHECK3-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, %class.anon.0*)* @__omp_outlined__2 to i8*), i8* null, i8** [[TMP5]], i64 1)
|
|
// CHECK3-NEXT: br label [[DOTOMP_DEINIT:%.*]]
|
|
// CHECK3: .omp.deinit:
|
|
// CHECK3-NEXT: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
|
|
// CHECK3-NEXT: br label [[DOTEXIT:%.*]]
|
|
// CHECK3: .exit:
|
|
// CHECK3-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__2
|
|
// CHECK3-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], %class.anon.0* nonnull align 8 dereferenceable(8) [[T:%.*]]) #[[ATTR1]] {
|
|
// CHECK3-NEXT: entry:
|
|
// CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK3-NEXT: [[T_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK3-NEXT: [[TMP:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK3-NEXT: [[T1:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8
|
|
// CHECK3-NEXT: [[_TMP2:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
|
|
// CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
|
|
// CHECK3-NEXT: store %class.anon.0* [[T]], %class.anon.0** [[T_ADDR]], align 8
|
|
// CHECK3-NEXT: [[TMP0:%.*]] = load %class.anon.0*, %class.anon.0** [[T_ADDR]], align 8
|
|
// CHECK3-NEXT: store %class.anon.0* [[TMP0]], %class.anon.0** [[TMP]], align 8
|
|
// CHECK3-NEXT: [[TMP1:%.*]] = load %class.anon.0*, %class.anon.0** [[TMP]], align 8
|
|
// CHECK3-NEXT: [[TMP2:%.*]] = bitcast %class.anon.0* [[T1]] to i8*
|
|
// CHECK3-NEXT: [[TMP3:%.*]] = bitcast %class.anon.0* [[TMP1]] to i8*
|
|
// CHECK3-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP2]], i8* align 8 [[TMP3]], i64 8, i1 false)
|
|
// CHECK3-NEXT: store %class.anon.0* [[T1]], %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK3-NEXT: [[TMP4:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK3-NEXT: [[TMP5:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK3-NEXT: [[CALL:%.*]] = call i32 @_ZZN1S3fooEvENKUlvE_clEv(%class.anon.0* nonnull dereferenceable(8) [[TMP5]]) #[[ATTR7]]
|
|
// CHECK3-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l41_worker
|
|
// CHECK4-SAME: () #[[ATTR0:[0-9]+]] {
|
|
// CHECK4-NEXT: entry:
|
|
// CHECK4-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 8
|
|
// CHECK4-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1
|
|
// CHECK4-NEXT: store i8* null, i8** [[WORK_FN]], align 8
|
|
// CHECK4-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1
|
|
// CHECK4-NEXT: br label [[DOTAWAIT_WORK:%.*]]
|
|
// CHECK4: .await.work:
|
|
// CHECK4-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
|
|
// CHECK4-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]])
|
|
// CHECK4-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8
|
|
// CHECK4-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1
|
|
// CHECK4-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 8
|
|
// CHECK4-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null
|
|
// CHECK4-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]]
|
|
// CHECK4: .select.workers:
|
|
// CHECK4-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1
|
|
// CHECK4-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0
|
|
// CHECK4-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]]
|
|
// CHECK4: .execute.parallel:
|
|
// CHECK4-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1:[0-9]+]])
|
|
// CHECK4-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)*
|
|
// CHECK4-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]])
|
|
// CHECK4-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]]
|
|
// CHECK4: .terminate.parallel:
|
|
// CHECK4-NEXT: call void @__kmpc_kernel_end_parallel()
|
|
// CHECK4-NEXT: br label [[DOTBARRIER_PARALLEL]]
|
|
// CHECK4: .barrier.parallel:
|
|
// CHECK4-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
|
|
// CHECK4-NEXT: br label [[DOTAWAIT_WORK]]
|
|
// CHECK4: .exit:
|
|
// CHECK4-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l41
|
|
// CHECK4-SAME: (i64 [[ARGC:%.*]], i32* nonnull align 4 dereferenceable(4) [[B:%.*]], i32* nonnull align 4 dereferenceable(4) [[C:%.*]], i32* [[D:%.*]], i32* nonnull align 4 dereferenceable(4) [[A:%.*]], %class.anon* nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR1:[0-9]+]] {
|
|
// CHECK4-NEXT: entry:
|
|
// CHECK4-NEXT: [[ARGC_ADDR:%.*]] = alloca i64, align 8
|
|
// CHECK4-NEXT: [[B_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[D_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[A_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[L_ADDR:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK4-NEXT: [[TMP:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[_TMP1:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[_TMP2:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK4-NEXT: [[L9:%.*]] = alloca [[CLASS_ANON:%.*]], align 8
|
|
// CHECK4-NEXT: [[_TMP10:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK4-NEXT: [[B11:%.*]] = alloca i32, align 4
|
|
// CHECK4-NEXT: [[_TMP12:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[C13:%.*]] = alloca i32, align 4
|
|
// CHECK4-NEXT: [[_TMP14:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: store i64 [[ARGC]], i64* [[ARGC_ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[B]], i32** [[B_ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[D]], i32** [[D_ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[A]], i32** [[A_ADDR]], align 8
|
|
// CHECK4-NEXT: store %class.anon* [[L]], %class.anon** [[L_ADDR]], align 8
|
|
// CHECK4-NEXT: [[CONV:%.*]] = bitcast i64* [[ARGC_ADDR]] to i32*
|
|
// CHECK4-NEXT: [[TMP0:%.*]] = load i32*, i32** [[B_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP1:%.*]] = load i32*, i32** [[C_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP2:%.*]] = load i32*, i32** [[A_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP3:%.*]] = load %class.anon*, %class.anon** [[L_ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[TMP0]], i32** [[TMP]], align 8
|
|
// CHECK4-NEXT: store i32* [[TMP1]], i32** [[_TMP1]], align 8
|
|
// CHECK4-NEXT: store %class.anon* [[TMP3]], %class.anon** [[_TMP2]], align 8
|
|
// CHECK4-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
|
|
// CHECK4-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK4-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
|
|
// CHECK4-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]]
|
|
// CHECK4-NEXT: [[TMP4:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]]
|
|
// CHECK4-NEXT: br i1 [[TMP4]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]]
|
|
// CHECK4: .worker:
|
|
// CHECK4-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l41_worker() #[[ATTR6:[0-9]+]]
|
|
// CHECK4-NEXT: br label [[DOTEXIT:%.*]]
|
|
// CHECK4: .mastercheck:
|
|
// CHECK4-NEXT: [[NVPTX_TID3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
|
|
// CHECK4-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK4-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
|
|
// CHECK4-NEXT: [[TMP5:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE5]], 1
|
|
// CHECK4-NEXT: [[TMP6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], 1
|
|
// CHECK4-NEXT: [[TMP7:%.*]] = xor i32 [[TMP5]], -1
|
|
// CHECK4-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP6]], [[TMP7]]
|
|
// CHECK4-NEXT: [[TMP8:%.*]] = icmp eq i32 [[NVPTX_TID3]], [[MASTER_TID]]
|
|
// CHECK4-NEXT: br i1 [[TMP8]], label [[DOTMASTER:%.*]], label [[DOTEXIT]]
|
|
// CHECK4: .master:
|
|
// CHECK4-NEXT: [[NVPTX_NUM_THREADS6:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK4-NEXT: [[NVPTX_WARP_SIZE7:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
|
|
// CHECK4-NEXT: [[THREAD_LIMIT8:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS6]], [[NVPTX_WARP_SIZE7]]
|
|
// CHECK4-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT8]], i16 1)
|
|
// CHECK4-NEXT: call void @__kmpc_data_sharing_init_stack()
|
|
// CHECK4-NEXT: [[TMP9:%.*]] = load %class.anon*, %class.anon** [[_TMP2]], align 8
|
|
// CHECK4-NEXT: [[TMP10:%.*]] = bitcast %class.anon* [[L9]] to i8*
|
|
// CHECK4-NEXT: [[TMP11:%.*]] = bitcast %class.anon* [[TMP9]] to i8*
|
|
// CHECK4-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP10]], i8* align 8 [[TMP11]], i64 40, i1 false)
|
|
// CHECK4-NEXT: store %class.anon* [[L9]], %class.anon** [[_TMP10]], align 8
|
|
// CHECK4-NEXT: [[TMP12:%.*]] = load i32*, i32** [[TMP]], align 8
|
|
// CHECK4-NEXT: [[TMP13:%.*]] = load i32, i32* [[TMP12]], align 4
|
|
// CHECK4-NEXT: store i32 [[TMP13]], i32* [[B11]], align 4
|
|
// CHECK4-NEXT: store i32* [[B11]], i32** [[_TMP12]], align 8
|
|
// CHECK4-NEXT: [[TMP14:%.*]] = load i32*, i32** [[_TMP1]], align 8
|
|
// CHECK4-NEXT: [[TMP15:%.*]] = load i32, i32* [[TMP14]], align 4
|
|
// CHECK4-NEXT: store i32 [[TMP15]], i32* [[C13]], align 4
|
|
// CHECK4-NEXT: store i32* [[C13]], i32** [[_TMP14]], align 8
|
|
// CHECK4-NEXT: [[TMP16:%.*]] = load %class.anon*, %class.anon** [[_TMP10]], align 8
|
|
// CHECK4-NEXT: [[TMP17:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP16]], i32 0, i32 0
|
|
// CHECK4-NEXT: store i32* [[CONV]], i32** [[TMP17]], align 8
|
|
// CHECK4-NEXT: [[TMP18:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP16]], i32 0, i32 1
|
|
// CHECK4-NEXT: [[TMP19:%.*]] = load i32*, i32** [[_TMP12]], align 8
|
|
// CHECK4-NEXT: store i32* [[TMP19]], i32** [[TMP18]], align 8
|
|
// CHECK4-NEXT: [[TMP20:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP16]], i32 0, i32 2
|
|
// CHECK4-NEXT: [[TMP21:%.*]] = load i32*, i32** [[_TMP14]], align 8
|
|
// CHECK4-NEXT: store i32* [[TMP21]], i32** [[TMP20]], align 8
|
|
// CHECK4-NEXT: [[TMP22:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP16]], i32 0, i32 3
|
|
// CHECK4-NEXT: store i32** [[D_ADDR]], i32*** [[TMP22]], align 8
|
|
// CHECK4-NEXT: [[TMP23:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP16]], i32 0, i32 4
|
|
// CHECK4-NEXT: store i32* [[TMP2]], i32** [[TMP23]], align 8
|
|
// CHECK4-NEXT: [[TMP24:%.*]] = load %class.anon*, %class.anon** [[_TMP10]], align 8
|
|
// CHECK4-NEXT: [[CALL:%.*]] = call i64 @"_ZZ4mainENK3$_0clEv"(%class.anon* nonnull dereferenceable(40) [[TMP24]]) #[[ATTR7:[0-9]+]]
|
|
// CHECK4-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]]
|
|
// CHECK4: .termination.notifier:
|
|
// CHECK4-NEXT: call void @__kmpc_kernel_deinit(i16 1)
|
|
// CHECK4-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
|
|
// CHECK4-NEXT: br label [[DOTEXIT]]
|
|
// CHECK4: .exit:
|
|
// CHECK4-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l43
|
|
// CHECK4-SAME: (i32* nonnull align 4 dereferenceable(4) [[ARGC:%.*]], i32* nonnull align 4 dereferenceable(4) [[B:%.*]], i32* nonnull align 4 dereferenceable(4) [[C:%.*]], i32* [[D:%.*]], i32* nonnull align 4 dereferenceable(4) [[A:%.*]], %class.anon* nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK4-NEXT: entry:
|
|
// CHECK4-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[B_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[D_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[A_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[L_ADDR:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK4-NEXT: [[TMP:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[_TMP1:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[_TMP2:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK4-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [6 x i8*], align 8
|
|
// CHECK4-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[B]], i32** [[B_ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[D]], i32** [[D_ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[A]], i32** [[A_ADDR]], align 8
|
|
// CHECK4-NEXT: store %class.anon* [[L]], %class.anon** [[L_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP1:%.*]] = load i32*, i32** [[B_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP2:%.*]] = load i32*, i32** [[C_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP3:%.*]] = load i32*, i32** [[A_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP4:%.*]] = load %class.anon*, %class.anon** [[L_ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[TMP1]], i32** [[TMP]], align 8
|
|
// CHECK4-NEXT: store i32* [[TMP2]], i32** [[_TMP1]], align 8
|
|
// CHECK4-NEXT: store %class.anon* [[TMP4]], %class.anon** [[_TMP2]], align 8
|
|
// CHECK4-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK4-NEXT: call void @__kmpc_spmd_kernel_init(i32 [[NVPTX_NUM_THREADS]], i16 1)
|
|
// CHECK4-NEXT: call void @__kmpc_data_sharing_init_stack_spmd()
|
|
// CHECK4-NEXT: br label [[DOTEXECUTE:%.*]]
|
|
// CHECK4: .execute:
|
|
// CHECK4-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2:[0-9]+]])
|
|
// CHECK4-NEXT: [[TMP6:%.*]] = load i32*, i32** [[TMP]], align 8
|
|
// CHECK4-NEXT: [[TMP7:%.*]] = load i32*, i32** [[_TMP1]], align 8
|
|
// CHECK4-NEXT: [[TMP8:%.*]] = load i32*, i32** [[D_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP9:%.*]] = load %class.anon*, %class.anon** [[_TMP2]], align 8
|
|
// CHECK4-NEXT: [[TMP10:%.*]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0
|
|
// CHECK4-NEXT: [[TMP11:%.*]] = bitcast i32* [[TMP0]] to i8*
|
|
// CHECK4-NEXT: store i8* [[TMP11]], i8** [[TMP10]], align 8
|
|
// CHECK4-NEXT: [[TMP12:%.*]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 1
|
|
// CHECK4-NEXT: [[TMP13:%.*]] = bitcast i32* [[TMP6]] to i8*
|
|
// CHECK4-NEXT: store i8* [[TMP13]], i8** [[TMP12]], align 8
|
|
// CHECK4-NEXT: [[TMP14:%.*]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 2
|
|
// CHECK4-NEXT: [[TMP15:%.*]] = bitcast i32* [[TMP7]] to i8*
|
|
// CHECK4-NEXT: store i8* [[TMP15]], i8** [[TMP14]], align 8
|
|
// CHECK4-NEXT: [[TMP16:%.*]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 3
|
|
// CHECK4-NEXT: [[TMP17:%.*]] = bitcast i32* [[TMP8]] to i8*
|
|
// CHECK4-NEXT: store i8* [[TMP17]], i8** [[TMP16]], align 8
|
|
// CHECK4-NEXT: [[TMP18:%.*]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 4
|
|
// CHECK4-NEXT: [[TMP19:%.*]] = bitcast i32* [[TMP3]] to i8*
|
|
// CHECK4-NEXT: store i8* [[TMP19]], i8** [[TMP18]], align 8
|
|
// CHECK4-NEXT: [[TMP20:%.*]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 5
|
|
// CHECK4-NEXT: [[TMP21:%.*]] = bitcast %class.anon* [[TMP9]] to i8*
|
|
// CHECK4-NEXT: store i8* [[TMP21]], i8** [[TMP20]], align 8
|
|
// CHECK4-NEXT: [[TMP22:%.*]] = bitcast [6 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
|
|
// CHECK4-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP5]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32*, i32*, i32*, i32*, i32*, %class.anon*)* @__omp_outlined__ to i8*), i8* null, i8** [[TMP22]], i64 6)
|
|
// CHECK4-NEXT: br label [[DOTOMP_DEINIT:%.*]]
|
|
// CHECK4: .omp.deinit:
|
|
// CHECK4-NEXT: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
|
|
// CHECK4-NEXT: br label [[DOTEXIT:%.*]]
|
|
// CHECK4: .exit:
|
|
// CHECK4-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK4-LABEL: define {{[^@]+}}@__omp_outlined__
|
|
// CHECK4-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[ARGC:%.*]], i32* nonnull align 4 dereferenceable(4) [[B:%.*]], i32* nonnull align 4 dereferenceable(4) [[C:%.*]], i32* [[D:%.*]], i32* nonnull align 4 dereferenceable(4) [[A:%.*]], %class.anon* nonnull align 8 dereferenceable(40) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK4-NEXT: entry:
|
|
// CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[ARGC_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[B_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[C_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[D_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[A_ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[L_ADDR:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK4-NEXT: [[TMP:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[_TMP1:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[_TMP2:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK4-NEXT: [[L3:%.*]] = alloca [[CLASS_ANON:%.*]], align 8
|
|
// CHECK4-NEXT: [[_TMP4:%.*]] = alloca %class.anon*, align 8
|
|
// CHECK4-NEXT: [[ARGC5:%.*]] = alloca i32, align 4
|
|
// CHECK4-NEXT: [[B6:%.*]] = alloca i32, align 4
|
|
// CHECK4-NEXT: [[_TMP7:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[C8:%.*]] = alloca i32, align 4
|
|
// CHECK4-NEXT: [[_TMP9:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[A10:%.*]] = alloca i32, align 4
|
|
// CHECK4-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[ARGC]], i32** [[ARGC_ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[B]], i32** [[B_ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[C]], i32** [[C_ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[D]], i32** [[D_ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[A]], i32** [[A_ADDR]], align 8
|
|
// CHECK4-NEXT: store %class.anon* [[L]], %class.anon** [[L_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP0:%.*]] = load i32*, i32** [[ARGC_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP1:%.*]] = load i32*, i32** [[B_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP2:%.*]] = load i32*, i32** [[C_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP3:%.*]] = load i32*, i32** [[A_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP4:%.*]] = load %class.anon*, %class.anon** [[L_ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[TMP1]], i32** [[TMP]], align 8
|
|
// CHECK4-NEXT: store i32* [[TMP2]], i32** [[_TMP1]], align 8
|
|
// CHECK4-NEXT: store %class.anon* [[TMP4]], %class.anon** [[_TMP2]], align 8
|
|
// CHECK4-NEXT: [[TMP5:%.*]] = load %class.anon*, %class.anon** [[_TMP2]], align 8
|
|
// CHECK4-NEXT: [[TMP6:%.*]] = bitcast %class.anon* [[L3]] to i8*
|
|
// CHECK4-NEXT: [[TMP7:%.*]] = bitcast %class.anon* [[TMP5]] to i8*
|
|
// CHECK4-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP6]], i8* align 8 [[TMP7]], i64 40, i1 false)
|
|
// CHECK4-NEXT: store %class.anon* [[L3]], %class.anon** [[_TMP4]], align 8
|
|
// CHECK4-NEXT: [[TMP8:%.*]] = load i32, i32* [[TMP0]], align 4
|
|
// CHECK4-NEXT: store i32 [[TMP8]], i32* [[ARGC5]], align 4
|
|
// CHECK4-NEXT: [[TMP9:%.*]] = load i32*, i32** [[TMP]], align 8
|
|
// CHECK4-NEXT: [[TMP10:%.*]] = load i32, i32* [[TMP9]], align 4
|
|
// CHECK4-NEXT: store i32 [[TMP10]], i32* [[B6]], align 4
|
|
// CHECK4-NEXT: store i32* [[B6]], i32** [[_TMP7]], align 8
|
|
// CHECK4-NEXT: [[TMP11:%.*]] = load i32*, i32** [[_TMP1]], align 8
|
|
// CHECK4-NEXT: [[TMP12:%.*]] = load i32, i32* [[TMP11]], align 4
|
|
// CHECK4-NEXT: store i32 [[TMP12]], i32* [[C8]], align 4
|
|
// CHECK4-NEXT: store i32* [[C8]], i32** [[_TMP9]], align 8
|
|
// CHECK4-NEXT: [[TMP13:%.*]] = load i32, i32* [[TMP3]], align 4
|
|
// CHECK4-NEXT: store i32 [[TMP13]], i32* [[A10]], align 4
|
|
// CHECK4-NEXT: [[TMP14:%.*]] = load %class.anon*, %class.anon** [[_TMP4]], align 8
|
|
// CHECK4-NEXT: [[TMP15:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP14]], i32 0, i32 0
|
|
// CHECK4-NEXT: store i32* [[ARGC5]], i32** [[TMP15]], align 8
|
|
// CHECK4-NEXT: [[TMP16:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP14]], i32 0, i32 1
|
|
// CHECK4-NEXT: [[TMP17:%.*]] = load i32*, i32** [[_TMP7]], align 8
|
|
// CHECK4-NEXT: store i32* [[TMP17]], i32** [[TMP16]], align 8
|
|
// CHECK4-NEXT: [[TMP18:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP14]], i32 0, i32 2
|
|
// CHECK4-NEXT: [[TMP19:%.*]] = load i32*, i32** [[_TMP9]], align 8
|
|
// CHECK4-NEXT: store i32* [[TMP19]], i32** [[TMP18]], align 8
|
|
// CHECK4-NEXT: [[TMP20:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP14]], i32 0, i32 3
|
|
// CHECK4-NEXT: store i32** [[D_ADDR]], i32*** [[TMP20]], align 8
|
|
// CHECK4-NEXT: [[TMP21:%.*]] = getelementptr inbounds [[CLASS_ANON]], %class.anon* [[TMP14]], i32 0, i32 4
|
|
// CHECK4-NEXT: store i32* [[A10]], i32** [[TMP21]], align 8
|
|
// CHECK4-NEXT: [[TMP22:%.*]] = load %class.anon*, %class.anon** [[_TMP4]], align 8
|
|
// CHECK4-NEXT: [[CALL:%.*]] = call i64 @"_ZZ4mainENK3$_0clEv"(%class.anon* nonnull dereferenceable(40) [[TMP22]]) #[[ATTR7]]
|
|
// CHECK4-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l27_worker
|
|
// CHECK4-SAME: () #[[ATTR0]] {
|
|
// CHECK4-NEXT: entry:
|
|
// CHECK4-NEXT: [[WORK_FN:%.*]] = alloca i8*, align 8
|
|
// CHECK4-NEXT: [[EXEC_STATUS:%.*]] = alloca i8, align 1
|
|
// CHECK4-NEXT: store i8* null, i8** [[WORK_FN]], align 8
|
|
// CHECK4-NEXT: store i8 0, i8* [[EXEC_STATUS]], align 1
|
|
// CHECK4-NEXT: br label [[DOTAWAIT_WORK:%.*]]
|
|
// CHECK4: .await.work:
|
|
// CHECK4-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
|
|
// CHECK4-NEXT: [[TMP0:%.*]] = call i1 @__kmpc_kernel_parallel(i8** [[WORK_FN]])
|
|
// CHECK4-NEXT: [[TMP1:%.*]] = zext i1 [[TMP0]] to i8
|
|
// CHECK4-NEXT: store i8 [[TMP1]], i8* [[EXEC_STATUS]], align 1
|
|
// CHECK4-NEXT: [[TMP2:%.*]] = load i8*, i8** [[WORK_FN]], align 8
|
|
// CHECK4-NEXT: [[SHOULD_TERMINATE:%.*]] = icmp eq i8* [[TMP2]], null
|
|
// CHECK4-NEXT: br i1 [[SHOULD_TERMINATE]], label [[DOTEXIT:%.*]], label [[DOTSELECT_WORKERS:%.*]]
|
|
// CHECK4: .select.workers:
|
|
// CHECK4-NEXT: [[TMP3:%.*]] = load i8, i8* [[EXEC_STATUS]], align 1
|
|
// CHECK4-NEXT: [[IS_ACTIVE:%.*]] = icmp ne i8 [[TMP3]], 0
|
|
// CHECK4-NEXT: br i1 [[IS_ACTIVE]], label [[DOTEXECUTE_PARALLEL:%.*]], label [[DOTBARRIER_PARALLEL:%.*]]
|
|
// CHECK4: .execute.parallel:
|
|
// CHECK4-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB1]])
|
|
// CHECK4-NEXT: [[TMP5:%.*]] = bitcast i8* [[TMP2]] to void (i16, i32)*
|
|
// CHECK4-NEXT: call void [[TMP5]](i16 0, i32 [[TMP4]])
|
|
// CHECK4-NEXT: br label [[DOTTERMINATE_PARALLEL:%.*]]
|
|
// CHECK4: .terminate.parallel:
|
|
// CHECK4-NEXT: call void @__kmpc_kernel_end_parallel()
|
|
// CHECK4-NEXT: br label [[DOTBARRIER_PARALLEL]]
|
|
// CHECK4: .barrier.parallel:
|
|
// CHECK4-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
|
|
// CHECK4-NEXT: br label [[DOTAWAIT_WORK]]
|
|
// CHECK4: .exit:
|
|
// CHECK4-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l27
|
|
// CHECK4-SAME: (%struct.S* [[THIS:%.*]], %class.anon.0* nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK4-NEXT: entry:
|
|
// CHECK4-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.S*, align 8
|
|
// CHECK4-NEXT: [[L_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK4-NEXT: [[TMP:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK4-NEXT: [[L7:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8
|
|
// CHECK4-NEXT: [[_TMP8:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK4-NEXT: store %struct.S* [[THIS]], %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK4-NEXT: store %class.anon.0* [[L]], %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP0:%.*]] = load %struct.S*, %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP1:%.*]] = load %class.anon.0*, %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK4-NEXT: store %class.anon.0* [[TMP1]], %class.anon.0** [[TMP]], align 8
|
|
// CHECK4-NEXT: [[NVPTX_TID:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
|
|
// CHECK4-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK4-NEXT: [[NVPTX_WARP_SIZE:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
|
|
// CHECK4-NEXT: [[THREAD_LIMIT:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS]], [[NVPTX_WARP_SIZE]]
|
|
// CHECK4-NEXT: [[TMP2:%.*]] = icmp ult i32 [[NVPTX_TID]], [[THREAD_LIMIT]]
|
|
// CHECK4-NEXT: br i1 [[TMP2]], label [[DOTWORKER:%.*]], label [[DOTMASTERCHECK:%.*]]
|
|
// CHECK4: .worker:
|
|
// CHECK4-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l27_worker() #[[ATTR6]]
|
|
// CHECK4-NEXT: br label [[DOTEXIT:%.*]]
|
|
// CHECK4: .mastercheck:
|
|
// CHECK4-NEXT: [[NVPTX_TID1:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
|
|
// CHECK4-NEXT: [[NVPTX_NUM_THREADS2:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK4-NEXT: [[NVPTX_WARP_SIZE3:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
|
|
// CHECK4-NEXT: [[TMP3:%.*]] = sub nuw i32 [[NVPTX_WARP_SIZE3]], 1
|
|
// CHECK4-NEXT: [[TMP4:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS2]], 1
|
|
// CHECK4-NEXT: [[TMP5:%.*]] = xor i32 [[TMP3]], -1
|
|
// CHECK4-NEXT: [[MASTER_TID:%.*]] = and i32 [[TMP4]], [[TMP5]]
|
|
// CHECK4-NEXT: [[TMP6:%.*]] = icmp eq i32 [[NVPTX_TID1]], [[MASTER_TID]]
|
|
// CHECK4-NEXT: br i1 [[TMP6]], label [[DOTMASTER:%.*]], label [[DOTEXIT]]
|
|
// CHECK4: .master:
|
|
// CHECK4-NEXT: [[NVPTX_NUM_THREADS4:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK4-NEXT: [[NVPTX_WARP_SIZE5:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
|
|
// CHECK4-NEXT: [[THREAD_LIMIT6:%.*]] = sub nuw i32 [[NVPTX_NUM_THREADS4]], [[NVPTX_WARP_SIZE5]]
|
|
// CHECK4-NEXT: call void @__kmpc_kernel_init(i32 [[THREAD_LIMIT6]], i16 1)
|
|
// CHECK4-NEXT: call void @__kmpc_data_sharing_init_stack()
|
|
// CHECK4-NEXT: [[TMP7:%.*]] = load %class.anon.0*, %class.anon.0** [[TMP]], align 8
|
|
// CHECK4-NEXT: [[TMP8:%.*]] = bitcast %class.anon.0* [[L7]] to i8*
|
|
// CHECK4-NEXT: [[TMP9:%.*]] = bitcast %class.anon.0* [[TMP7]] to i8*
|
|
// CHECK4-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP8]], i8* align 8 [[TMP9]], i64 8, i1 false)
|
|
// CHECK4-NEXT: store %class.anon.0* [[L7]], %class.anon.0** [[_TMP8]], align 8
|
|
// CHECK4-NEXT: [[TMP10:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP8]], align 8
|
|
// CHECK4-NEXT: [[TMP11:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP10]], i32 0, i32 0
|
|
// CHECK4-NEXT: store %struct.S* [[TMP0]], %struct.S** [[TMP11]], align 8
|
|
// CHECK4-NEXT: [[TMP12:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP8]], align 8
|
|
// CHECK4-NEXT: [[CALL:%.*]] = call i32 @_ZZN1S3fooEvENKUlvE_clEv(%class.anon.0* nonnull dereferenceable(8) [[TMP12]]) #[[ATTR7]]
|
|
// CHECK4-NEXT: br label [[DOTTERMINATION_NOTIFIER:%.*]]
|
|
// CHECK4: .termination.notifier:
|
|
// CHECK4-NEXT: call void @__kmpc_kernel_deinit(i16 1)
|
|
// CHECK4-NEXT: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
|
|
// CHECK4-NEXT: br label [[DOTEXIT]]
|
|
// CHECK4: .exit:
|
|
// CHECK4-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK4-LABEL: define {{[^@]+}}@_ZZN1S3fooEvENKUlvE_clEv
|
|
// CHECK4-SAME: (%class.anon.0* nonnull dereferenceable(8) [[THIS:%.*]]) #[[ATTR4:[0-9]+]] comdat align 2 {
|
|
// CHECK4-NEXT: entry:
|
|
// CHECK4-NEXT: [[THIS_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK4-NEXT: store %class.anon.0* [[THIS]], %class.anon.0** [[THIS_ADDR]], align 8
|
|
// CHECK4-NEXT: [[THIS1:%.*]] = load %class.anon.0*, %class.anon.0** [[THIS_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP0:%.*]] = getelementptr inbounds [[CLASS_ANON_0:%.*]], %class.anon.0* [[THIS1]], i32 0, i32 0
|
|
// CHECK4-NEXT: [[TMP1:%.*]] = load %struct.S*, %struct.S** [[TMP0]], align 8
|
|
// CHECK4-NEXT: [[A:%.*]] = getelementptr inbounds [[STRUCT_S:%.*]], %struct.S* [[TMP1]], i32 0, i32 0
|
|
// CHECK4-NEXT: [[TMP2:%.*]] = load i32, i32* [[A]], align 4
|
|
// CHECK4-NEXT: ret i32 [[TMP2]]
|
|
//
|
|
//
|
|
// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1S3fooEv_l29
|
|
// CHECK4-SAME: (%struct.S* [[THIS:%.*]], %class.anon.0* nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK4-NEXT: entry:
|
|
// CHECK4-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.S*, align 8
|
|
// CHECK4-NEXT: [[L_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK4-NEXT: [[TMP:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK4-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [2 x i8*], align 8
|
|
// CHECK4-NEXT: store %struct.S* [[THIS]], %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK4-NEXT: store %class.anon.0* [[L]], %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP0:%.*]] = load %struct.S*, %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP1:%.*]] = load %class.anon.0*, %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK4-NEXT: store %class.anon.0* [[TMP1]], %class.anon.0** [[TMP]], align 8
|
|
// CHECK4-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK4-NEXT: call void @__kmpc_spmd_kernel_init(i32 [[NVPTX_NUM_THREADS]], i16 1)
|
|
// CHECK4-NEXT: call void @__kmpc_data_sharing_init_stack_spmd()
|
|
// CHECK4-NEXT: br label [[DOTEXECUTE:%.*]]
|
|
// CHECK4: .execute:
|
|
// CHECK4-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2]])
|
|
// CHECK4-NEXT: [[TMP3:%.*]] = load %class.anon.0*, %class.anon.0** [[TMP]], align 8
|
|
// CHECK4-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0
|
|
// CHECK4-NEXT: [[TMP5:%.*]] = bitcast %struct.S* [[TMP0]] to i8*
|
|
// CHECK4-NEXT: store i8* [[TMP5]], i8** [[TMP4]], align 8
|
|
// CHECK4-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 1
|
|
// CHECK4-NEXT: [[TMP7:%.*]] = bitcast %class.anon.0* [[TMP3]] to i8*
|
|
// CHECK4-NEXT: store i8* [[TMP7]], i8** [[TMP6]], align 8
|
|
// CHECK4-NEXT: [[TMP8:%.*]] = bitcast [2 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
|
|
// CHECK4-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP2]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, %struct.S*, %class.anon.0*)* @__omp_outlined__1 to i8*), i8* null, i8** [[TMP8]], i64 2)
|
|
// CHECK4-NEXT: br label [[DOTOMP_DEINIT:%.*]]
|
|
// CHECK4: .omp.deinit:
|
|
// CHECK4-NEXT: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
|
|
// CHECK4-NEXT: br label [[DOTEXIT:%.*]]
|
|
// CHECK4: .exit:
|
|
// CHECK4-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK4-LABEL: define {{[^@]+}}@__omp_outlined__1
|
|
// CHECK4-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], %struct.S* [[THIS:%.*]], %class.anon.0* nonnull align 8 dereferenceable(8) [[L:%.*]]) #[[ATTR1]] {
|
|
// CHECK4-NEXT: entry:
|
|
// CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.S*, align 8
|
|
// CHECK4-NEXT: [[L_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK4-NEXT: [[TMP:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK4-NEXT: [[L1:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8
|
|
// CHECK4-NEXT: [[_TMP2:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK4-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
|
|
// CHECK4-NEXT: store %struct.S* [[THIS]], %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK4-NEXT: store %class.anon.0* [[L]], %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP0:%.*]] = load %struct.S*, %struct.S** [[THIS_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP1:%.*]] = load %class.anon.0*, %class.anon.0** [[L_ADDR]], align 8
|
|
// CHECK4-NEXT: store %class.anon.0* [[TMP1]], %class.anon.0** [[TMP]], align 8
|
|
// CHECK4-NEXT: [[TMP2:%.*]] = load %class.anon.0*, %class.anon.0** [[TMP]], align 8
|
|
// CHECK4-NEXT: [[TMP3:%.*]] = bitcast %class.anon.0* [[L1]] to i8*
|
|
// CHECK4-NEXT: [[TMP4:%.*]] = bitcast %class.anon.0* [[TMP2]] to i8*
|
|
// CHECK4-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP3]], i8* align 8 [[TMP4]], i64 8, i1 false)
|
|
// CHECK4-NEXT: store %class.anon.0* [[L1]], %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK4-NEXT: [[TMP5:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK4-NEXT: [[TMP6:%.*]] = getelementptr inbounds [[CLASS_ANON_0]], %class.anon.0* [[TMP5]], i32 0, i32 0
|
|
// CHECK4-NEXT: store %struct.S* [[TMP0]], %struct.S** [[TMP6]], align 8
|
|
// CHECK4-NEXT: [[TMP7:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK4-NEXT: [[CALL:%.*]] = call i32 @_ZZN1S3fooEvENKUlvE_clEv(%class.anon.0* nonnull dereferenceable(8) [[TMP7]]) #[[ATTR7]]
|
|
// CHECK4-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooIZN1S3fooEvEUlvE_EiRKT__l18
|
|
// CHECK4-SAME: (%class.anon.0* nonnull align 8 dereferenceable(8) [[T:%.*]]) #[[ATTR1]] {
|
|
// CHECK4-NEXT: entry:
|
|
// CHECK4-NEXT: [[T_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK4-NEXT: [[TMP:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK4-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8
|
|
// CHECK4-NEXT: store %class.anon.0* [[T]], %class.anon.0** [[T_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP0:%.*]] = load %class.anon.0*, %class.anon.0** [[T_ADDR]], align 8
|
|
// CHECK4-NEXT: store %class.anon.0* [[TMP0]], %class.anon.0** [[TMP]], align 8
|
|
// CHECK4-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
|
// CHECK4-NEXT: call void @__kmpc_spmd_kernel_init(i32 [[NVPTX_NUM_THREADS]], i16 1)
|
|
// CHECK4-NEXT: call void @__kmpc_data_sharing_init_stack_spmd()
|
|
// CHECK4-NEXT: br label [[DOTEXECUTE:%.*]]
|
|
// CHECK4: .execute:
|
|
// CHECK4-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2]])
|
|
// CHECK4-NEXT: [[TMP2:%.*]] = load %class.anon.0*, %class.anon.0** [[TMP]], align 8
|
|
// CHECK4-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[CAPTURED_VARS_ADDRS]], i64 0, i64 0
|
|
// CHECK4-NEXT: [[TMP4:%.*]] = bitcast %class.anon.0* [[TMP2]] to i8*
|
|
// CHECK4-NEXT: store i8* [[TMP4]], i8** [[TMP3]], align 8
|
|
// CHECK4-NEXT: [[TMP5:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
|
|
// CHECK4-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, %class.anon.0*)* @__omp_outlined__2 to i8*), i8* null, i8** [[TMP5]], i64 1)
|
|
// CHECK4-NEXT: br label [[DOTOMP_DEINIT:%.*]]
|
|
// CHECK4: .omp.deinit:
|
|
// CHECK4-NEXT: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
|
|
// CHECK4-NEXT: br label [[DOTEXIT:%.*]]
|
|
// CHECK4: .exit:
|
|
// CHECK4-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK4-LABEL: define {{[^@]+}}@__omp_outlined__2
|
|
// CHECK4-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], %class.anon.0* nonnull align 8 dereferenceable(8) [[T:%.*]]) #[[ATTR1]] {
|
|
// CHECK4-NEXT: entry:
|
|
// CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
|
|
// CHECK4-NEXT: [[T_ADDR:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK4-NEXT: [[TMP:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK4-NEXT: [[T1:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 8
|
|
// CHECK4-NEXT: [[_TMP2:%.*]] = alloca %class.anon.0*, align 8
|
|
// CHECK4-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
|
|
// CHECK4-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
|
|
// CHECK4-NEXT: store %class.anon.0* [[T]], %class.anon.0** [[T_ADDR]], align 8
|
|
// CHECK4-NEXT: [[TMP0:%.*]] = load %class.anon.0*, %class.anon.0** [[T_ADDR]], align 8
|
|
// CHECK4-NEXT: store %class.anon.0* [[TMP0]], %class.anon.0** [[TMP]], align 8
|
|
// CHECK4-NEXT: [[TMP1:%.*]] = load %class.anon.0*, %class.anon.0** [[TMP]], align 8
|
|
// CHECK4-NEXT: [[TMP2:%.*]] = bitcast %class.anon.0* [[T1]] to i8*
|
|
// CHECK4-NEXT: [[TMP3:%.*]] = bitcast %class.anon.0* [[TMP1]] to i8*
|
|
// CHECK4-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[TMP2]], i8* align 8 [[TMP3]], i64 8, i1 false)
|
|
// CHECK4-NEXT: store %class.anon.0* [[T1]], %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK4-NEXT: [[TMP4:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK4-NEXT: [[TMP5:%.*]] = load %class.anon.0*, %class.anon.0** [[_TMP2]], align 8
|
|
// CHECK4-NEXT: [[CALL:%.*]] = call i32 @_ZZN1S3fooEvENKUlvE_clEv(%class.anon.0* nonnull dereferenceable(8) [[TMP5]]) #[[ATTR7]]
|
|
// CHECK4-NEXT: ret void
|
|
//
|