https://github.com/ddpagan created https://github.com/llvm/llvm-project/pull/187831
Currently, for device compilation, references to complex functions creal/cimag/conj generate library calls. These in turn cause undefined symbol references when linking because they are not supported on the device. However, because they are/can be emitted inline, the calls aren't necessary. The problem is that these functions, normally considered to be builtins and then later inlined by CodeGen, lose their builtin status when compiling for the device. They should remain builtins regardless so they are still emitted inline. Testing: New test - clang/test/OpenMP/amdgcn_device_complex_builtins.c check-all >From 09cf87fd021a62f51d8148f4476e3fcf752b1a3a Mon Sep 17 00:00:00 2001 From: Dave Pagan <[email protected]> Date: Tue, 17 Mar 2026 21:17:13 -0500 Subject: [PATCH] [OpenMP][AMDGCN] creal/cimag/conj should be builtins/inlined on device. Currently, for device compilation, references to complex functions creal/cimag/conj generate library calls. These in turn cause undefined symbol references when linking because they are not supported on the device. However, because they are/can be emitted inline, the calls aren't necessary. The problem is that these functions, normally considered to be builtins and then later inlined by CodeGen, lose their builtin status when compiling for the device. They should remain builtins regardless so they are still emitted inline. Testing: New test - clang/test/OpenMP/amdgcn_device_complex_builtins.c check-all --- clang/lib/AST/Decl.cpp | 12 +- .../OpenMP/amdgcn_device_complex_builtins.c | 371 ++++++++++++++++++ 2 files changed, 380 insertions(+), 3 deletions(-) create mode 100644 clang/test/OpenMP/amdgcn_device_complex_builtins.c diff --git a/clang/lib/AST/Decl.cpp b/clang/lib/AST/Decl.cpp index 37b00eeca539c..4651ae35e7624 100644 --- a/clang/lib/AST/Decl.cpp +++ b/clang/lib/AST/Decl.cpp @@ -3810,12 +3810,18 @@ unsigned FunctionDecl::getBuiltinID(bool ConsiderWrapperFunctions) const { return 0; // As AMDGCN implementation of OpenMP does not have a device-side standard - // library, none of the predefined library functions except printf and malloc - // should be treated as a builtin i.e. 0 should be returned for them. + // library, none of the predefined library functions except printf, malloc, + // and the complex functions that can be emitted inline (creal, cimag, + // conj and their float/double variants) should be treated as a builtin + // i.e. 0 should be returned for them. The long double variants are excluded + // because AMDGCN does not support x86 long double. if (Context.getTargetInfo().getTriple().isAMDGCN() && Context.getLangOpts().OpenMPIsTargetDevice && Context.BuiltinInfo.isPredefinedLibFunction(BuiltinID) && - !(BuiltinID == Builtin::BIprintf || BuiltinID == Builtin::BImalloc)) + !(BuiltinID == Builtin::BIprintf || BuiltinID == Builtin::BImalloc || + BuiltinID == Builtin::BIcreal || BuiltinID == Builtin::BIcrealf || + BuiltinID == Builtin::BIcimag || BuiltinID == Builtin::BIcimagf || + BuiltinID == Builtin::BIconj || BuiltinID == Builtin::BIconjf)) return 0; return BuiltinID; diff --git a/clang/test/OpenMP/amdgcn_device_complex_builtins.c b/clang/test/OpenMP/amdgcn_device_complex_builtins.c new file mode 100644 index 0000000000000..c76ff0132e3c0 --- /dev/null +++ b/clang/test/OpenMP/amdgcn_device_complex_builtins.c @@ -0,0 +1,371 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" --prefix-filecheck-ir-name _ --version 5 +// REQUIRES: amdgpu-registered-target +// Verify that creal, cimag, and conj (and their float/double variants) are +// treated as builtins and emitted inline on both host and AMDGCN OpenMP target +// devices, rather than generating library calls. +// Long double variants (creall, cimagl, conjl) are excluded because AMDGCN +// does not support x86 long double. +// RUN: %clang_cc1 -fopenmp -x c -triple x86_64-unknown-unknown \ +// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc +// RUN: %clang_cc1 -fopenmp -x c -triple x86_64-unknown-unknown \ +// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=HOST +// RUN: %clang_cc1 -fopenmp -x c -triple amdgcn-amd-amdhsa \ +// RUN: -aux-triple x86_64-unknown-unknown \ +// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s \ +// RUN: -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-host.bc \ +// RUN: -o - | FileCheck %s --check-prefix=DEVICE + +extern float crealf(float _Complex); +extern double creal(double _Complex); + +extern float cimagf(float _Complex); +extern double cimag(double _Complex); + +extern float _Complex conjf(float _Complex); +extern double _Complex conj(double _Complex); + +void test_complex_builtins( + float _Complex fc, double _Complex dc, + float *fp, double *dp, + float _Complex *fcp, double _Complex *dcp) { + *fp = crealf(fc); + *dp = creal(dc); + *fp = cimagf(fc); + *dp = cimag(dc); + *fcp = conjf(fc); + *dcp = conj(dc); + +#pragma omp target + { + *fp = crealf(fc); + *dp = creal(dc); + *fp = cimagf(fc); + *dp = cimag(dc); + *fcp = conjf(fc); + *dcp = conj(dc); + } +} +// HOST-LABEL: define dso_local void @test_complex_builtins( +// HOST-SAME: <2 x float> noundef [[FC_COERCE:%.*]], double noundef [[DC_COERCE0:%.*]], double noundef [[DC_COERCE1:%.*]], ptr noundef [[FP:%.*]], ptr noundef [[DP:%.*]], ptr noundef [[FCP:%.*]], ptr noundef [[DCP:%.*]]) #[[ATTR0:[0-9]+]] { +// HOST-NEXT: [[ENTRY:.*:]] +// HOST-NEXT: [[FC:%.*]] = alloca { float, float }, align 4 +// HOST-NEXT: [[DC:%.*]] = alloca { double, double }, align 8 +// HOST-NEXT: [[FP_ADDR:%.*]] = alloca ptr, align 8 +// HOST-NEXT: [[DP_ADDR:%.*]] = alloca ptr, align 8 +// HOST-NEXT: [[FCP_ADDR:%.*]] = alloca ptr, align 8 +// HOST-NEXT: [[DCP_ADDR:%.*]] = alloca ptr, align 8 +// HOST-NEXT: [[FC_CASTED:%.*]] = alloca i64, align 8 +// HOST-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [7 x ptr], align 8 +// HOST-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [7 x ptr], align 8 +// HOST-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [7 x ptr], align 8 +// HOST-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8 +// HOST-NEXT: store <2 x float> [[FC_COERCE]], ptr [[FC]], align 4 +// HOST-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC]], i32 0, i32 0 +// HOST-NEXT: store double [[DC_COERCE0]], ptr [[TMP0]], align 8 +// HOST-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC]], i32 0, i32 1 +// HOST-NEXT: store double [[DC_COERCE1]], ptr [[TMP1]], align 8 +// HOST-NEXT: store ptr [[FP]], ptr [[FP_ADDR]], align 8 +// HOST-NEXT: store ptr [[DP]], ptr [[DP_ADDR]], align 8 +// HOST-NEXT: store ptr [[FCP]], ptr [[FCP_ADDR]], align 8 +// HOST-NEXT: store ptr [[DCP]], ptr [[DCP_ADDR]], align 8 +// HOST-NEXT: [[FC_REALP:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[FC]], i32 0, i32 0 +// HOST-NEXT: [[FC_REAL:%.*]] = load float, ptr [[FC_REALP]], align 4 +// HOST-NEXT: [[FC_IMAGP:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[FC]], i32 0, i32 1 +// HOST-NEXT: [[FC_IMAG:%.*]] = load float, ptr [[FC_IMAGP]], align 4 +// HOST-NEXT: [[TMP2:%.*]] = load ptr, ptr [[FP_ADDR]], align 8 +// HOST-NEXT: store float [[FC_REAL]], ptr [[TMP2]], align 4 +// HOST-NEXT: [[DC_REALP:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC]], i32 0, i32 0 +// HOST-NEXT: [[DC_REAL:%.*]] = load double, ptr [[DC_REALP]], align 8 +// HOST-NEXT: [[DC_IMAGP:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC]], i32 0, i32 1 +// HOST-NEXT: [[DC_IMAG:%.*]] = load double, ptr [[DC_IMAGP]], align 8 +// HOST-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DP_ADDR]], align 8 +// HOST-NEXT: store double [[DC_REAL]], ptr [[TMP3]], align 8 +// HOST-NEXT: [[FC_REALP1:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[FC]], i32 0, i32 0 +// HOST-NEXT: [[FC_REAL2:%.*]] = load float, ptr [[FC_REALP1]], align 4 +// HOST-NEXT: [[FC_IMAGP3:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[FC]], i32 0, i32 1 +// HOST-NEXT: [[FC_IMAG4:%.*]] = load float, ptr [[FC_IMAGP3]], align 4 +// HOST-NEXT: [[TMP4:%.*]] = load ptr, ptr [[FP_ADDR]], align 8 +// HOST-NEXT: store float [[FC_IMAG4]], ptr [[TMP4]], align 4 +// HOST-NEXT: [[DC_REALP5:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC]], i32 0, i32 0 +// HOST-NEXT: [[DC_REAL6:%.*]] = load double, ptr [[DC_REALP5]], align 8 +// HOST-NEXT: [[DC_IMAGP7:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC]], i32 0, i32 1 +// HOST-NEXT: [[DC_IMAG8:%.*]] = load double, ptr [[DC_IMAGP7]], align 8 +// HOST-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DP_ADDR]], align 8 +// HOST-NEXT: store double [[DC_IMAG8]], ptr [[TMP5]], align 8 +// HOST-NEXT: [[FC_REALP9:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[FC]], i32 0, i32 0 +// HOST-NEXT: [[FC_REAL10:%.*]] = load float, ptr [[FC_REALP9]], align 4 +// HOST-NEXT: [[FC_IMAGP11:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[FC]], i32 0, i32 1 +// HOST-NEXT: [[FC_IMAG12:%.*]] = load float, ptr [[FC_IMAGP11]], align 4 +// HOST-NEXT: [[NEG:%.*]] = fneg float [[FC_IMAG12]] +// HOST-NEXT: [[TMP6:%.*]] = load ptr, ptr [[FCP_ADDR]], align 8 +// HOST-NEXT: [[DOTREALP:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[TMP6]], i32 0, i32 0 +// HOST-NEXT: [[DOTIMAGP:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[TMP6]], i32 0, i32 1 +// HOST-NEXT: store float [[FC_REAL10]], ptr [[DOTREALP]], align 4 +// HOST-NEXT: store float [[NEG]], ptr [[DOTIMAGP]], align 4 +// HOST-NEXT: [[DC_REALP13:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC]], i32 0, i32 0 +// HOST-NEXT: [[DC_REAL14:%.*]] = load double, ptr [[DC_REALP13]], align 8 +// HOST-NEXT: [[DC_IMAGP15:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC]], i32 0, i32 1 +// HOST-NEXT: [[DC_IMAG16:%.*]] = load double, ptr [[DC_IMAGP15]], align 8 +// HOST-NEXT: [[NEG17:%.*]] = fneg double [[DC_IMAG16]] +// HOST-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DCP_ADDR]], align 8 +// HOST-NEXT: [[DOTREALP18:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[TMP7]], i32 0, i32 0 +// HOST-NEXT: [[DOTIMAGP19:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[TMP7]], i32 0, i32 1 +// HOST-NEXT: store double [[DC_REAL14]], ptr [[DOTREALP18]], align 8 +// HOST-NEXT: store double [[NEG17]], ptr [[DOTIMAGP19]], align 8 +// HOST-NEXT: [[TMP8:%.*]] = load ptr, ptr [[FP_ADDR]], align 8 +// HOST-NEXT: [[TMP9:%.*]] = load { float, float }, ptr [[FC]], align 4 +// HOST-NEXT: store { float, float } [[TMP9]], ptr [[FC_CASTED]], align 4 +// HOST-NEXT: [[TMP10:%.*]] = load i64, ptr [[FC_CASTED]], align 8 +// HOST-NEXT: [[TMP11:%.*]] = load ptr, ptr [[DP_ADDR]], align 8 +// HOST-NEXT: [[TMP12:%.*]] = load ptr, ptr [[FCP_ADDR]], align 8 +// HOST-NEXT: [[TMP13:%.*]] = load ptr, ptr [[DCP_ADDR]], align 8 +// HOST-NEXT: [[TMP14:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// HOST-NEXT: store ptr [[TMP8]], ptr [[TMP14]], align 8 +// HOST-NEXT: [[TMP15:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// HOST-NEXT: store ptr [[TMP8]], ptr [[TMP15]], align 8 +// HOST-NEXT: [[TMP16:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 +// HOST-NEXT: store ptr null, ptr [[TMP16]], align 8 +// HOST-NEXT: [[TMP17:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// HOST-NEXT: store i64 [[TMP10]], ptr [[TMP17]], align 8 +// HOST-NEXT: [[TMP18:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// HOST-NEXT: store i64 [[TMP10]], ptr [[TMP18]], align 8 +// HOST-NEXT: [[TMP19:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// HOST-NEXT: store ptr null, ptr [[TMP19]], align 8 +// HOST-NEXT: [[TMP20:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 +// HOST-NEXT: store ptr [[TMP11]], ptr [[TMP20]], align 8 +// HOST-NEXT: [[TMP21:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2 +// HOST-NEXT: store ptr [[TMP11]], ptr [[TMP21]], align 8 +// HOST-NEXT: [[TMP22:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2 +// HOST-NEXT: store ptr null, ptr [[TMP22]], align 8 +// HOST-NEXT: [[TMP23:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3 +// HOST-NEXT: store ptr [[DC]], ptr [[TMP23]], align 8 +// HOST-NEXT: [[TMP24:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3 +// HOST-NEXT: store ptr [[DC]], ptr [[TMP24]], align 8 +// HOST-NEXT: [[TMP25:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3 +// HOST-NEXT: store ptr null, ptr [[TMP25]], align 8 +// HOST-NEXT: [[TMP26:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4 +// HOST-NEXT: store ptr [[TMP12]], ptr [[TMP26]], align 8 +// HOST-NEXT: [[TMP27:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 4 +// HOST-NEXT: store ptr [[TMP12]], ptr [[TMP27]], align 8 +// HOST-NEXT: [[TMP28:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 4 +// HOST-NEXT: store ptr null, ptr [[TMP28]], align 8 +// HOST-NEXT: [[TMP29:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5 +// HOST-NEXT: store ptr [[TMP13]], ptr [[TMP29]], align 8 +// HOST-NEXT: [[TMP30:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 5 +// HOST-NEXT: store ptr [[TMP13]], ptr [[TMP30]], align 8 +// HOST-NEXT: [[TMP31:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 5 +// HOST-NEXT: store ptr null, ptr [[TMP31]], align 8 +// HOST-NEXT: [[TMP32:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 6 +// HOST-NEXT: store ptr null, ptr [[TMP32]], align 8 +// HOST-NEXT: [[TMP33:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 6 +// HOST-NEXT: store ptr null, ptr [[TMP33]], align 8 +// HOST-NEXT: [[TMP34:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 6 +// HOST-NEXT: store ptr null, ptr [[TMP34]], align 8 +// HOST-NEXT: [[TMP35:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// HOST-NEXT: [[TMP36:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// HOST-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0 +// HOST-NEXT: store i32 4, ptr [[TMP37]], align 4 +// HOST-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1 +// HOST-NEXT: store i32 7, ptr [[TMP38]], align 4 +// HOST-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2 +// HOST-NEXT: store ptr [[TMP35]], ptr [[TMP39]], align 8 +// HOST-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3 +// HOST-NEXT: store ptr [[TMP36]], ptr [[TMP40]], align 8 +// HOST-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4 +// HOST-NEXT: store ptr @.offload_sizes, ptr [[TMP41]], align 8 +// HOST-NEXT: [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5 +// HOST-NEXT: store ptr @.offload_maptypes, ptr [[TMP42]], align 8 +// HOST-NEXT: [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6 +// HOST-NEXT: store ptr null, ptr [[TMP43]], align 8 +// HOST-NEXT: [[TMP44:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7 +// HOST-NEXT: store ptr null, ptr [[TMP44]], align 8 +// HOST-NEXT: [[TMP45:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8 +// HOST-NEXT: store i64 0, ptr [[TMP45]], align 8 +// HOST-NEXT: [[TMP46:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9 +// HOST-NEXT: store i64 0, ptr [[TMP46]], align 8 +// HOST-NEXT: [[TMP47:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 +// HOST-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP47]], align 4 +// HOST-NEXT: [[TMP48:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 +// HOST-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP48]], align 4 +// HOST-NEXT: [[TMP49:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 +// HOST-NEXT: store i32 0, ptr [[TMP49]], align 4 +// HOST-NEXT: [[TMP50:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_test_complex_builtins_l39.region_id, ptr [[KERNEL_ARGS]]) +// HOST-NEXT: [[TMP51:%.*]] = icmp ne i32 [[TMP50]], 0 +// HOST-NEXT: br i1 [[TMP51]], label %[[OMP_OFFLOAD_FAILED:.*]], label %[[OMP_OFFLOAD_CONT:.*]] +// HOST: [[OMP_OFFLOAD_FAILED]]: +// HOST-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_test_complex_builtins_l39(ptr [[TMP8]], i64 [[TMP10]], ptr [[TMP11]], ptr [[DC]], ptr [[TMP12]], ptr [[TMP13]], ptr null) #[[ATTR2:[0-9]+]] +// HOST-NEXT: br label %[[OMP_OFFLOAD_CONT]] +// HOST: [[OMP_OFFLOAD_CONT]]: +// HOST-NEXT: ret void +// +// +// HOST-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_test_complex_builtins_l39( +// HOST-SAME: ptr noundef [[FP:%.*]], i64 noundef [[FC:%.*]], ptr noundef [[DP:%.*]], ptr noundef nonnull align 8 dereferenceable(16) [[DC:%.*]], ptr noundef [[FCP:%.*]], ptr noundef [[DCP:%.*]], ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR1:[0-9]+]] { +// HOST-NEXT: [[ENTRY:.*:]] +// HOST-NEXT: [[FP_ADDR:%.*]] = alloca ptr, align 8 +// HOST-NEXT: [[FC_ADDR:%.*]] = alloca i64, align 8 +// HOST-NEXT: [[DP_ADDR:%.*]] = alloca ptr, align 8 +// HOST-NEXT: [[DC_ADDR:%.*]] = alloca ptr, align 8 +// HOST-NEXT: [[FCP_ADDR:%.*]] = alloca ptr, align 8 +// HOST-NEXT: [[DCP_ADDR:%.*]] = alloca ptr, align 8 +// HOST-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8 +// HOST-NEXT: [[DC1:%.*]] = alloca { double, double }, align 8 +// HOST-NEXT: store ptr [[FP]], ptr [[FP_ADDR]], align 8 +// HOST-NEXT: store i64 [[FC]], ptr [[FC_ADDR]], align 8 +// HOST-NEXT: store ptr [[DP]], ptr [[DP_ADDR]], align 8 +// HOST-NEXT: store ptr [[DC]], ptr [[DC_ADDR]], align 8 +// HOST-NEXT: store ptr [[FCP]], ptr [[FCP_ADDR]], align 8 +// HOST-NEXT: store ptr [[DCP]], ptr [[DCP_ADDR]], align 8 +// HOST-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8 +// HOST-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DC_ADDR]], align 8, !nonnull [[META4:![0-9]+]], !align [[META5:![0-9]+]] +// HOST-NEXT: [[DOTREALP:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[TMP0]], i32 0, i32 0 +// HOST-NEXT: [[DOTREAL:%.*]] = load double, ptr [[DOTREALP]], align 8 +// HOST-NEXT: [[DOTIMAGP:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[TMP0]], i32 0, i32 1 +// HOST-NEXT: [[DOTIMAG:%.*]] = load double, ptr [[DOTIMAGP]], align 8 +// HOST-NEXT: [[DC1_REALP:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC1]], i32 0, i32 0 +// HOST-NEXT: [[DC1_IMAGP:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC1]], i32 0, i32 1 +// HOST-NEXT: store double [[DOTREAL]], ptr [[DC1_REALP]], align 8 +// HOST-NEXT: store double [[DOTIMAG]], ptr [[DC1_IMAGP]], align 8 +// HOST-NEXT: [[FC_ADDR_REALP:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[FC_ADDR]], i32 0, i32 0 +// HOST-NEXT: [[FC_ADDR_REAL:%.*]] = load float, ptr [[FC_ADDR_REALP]], align 4 +// HOST-NEXT: [[FC_ADDR_IMAGP:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[FC_ADDR]], i32 0, i32 1 +// HOST-NEXT: [[FC_ADDR_IMAG:%.*]] = load float, ptr [[FC_ADDR_IMAGP]], align 4 +// HOST-NEXT: [[TMP1:%.*]] = load ptr, ptr [[FP_ADDR]], align 8 +// HOST-NEXT: store float [[FC_ADDR_REAL]], ptr [[TMP1]], align 4 +// HOST-NEXT: [[DC1_REALP2:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC1]], i32 0, i32 0 +// HOST-NEXT: [[DC1_REAL:%.*]] = load double, ptr [[DC1_REALP2]], align 8 +// HOST-NEXT: [[DC1_IMAGP3:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC1]], i32 0, i32 1 +// HOST-NEXT: [[DC1_IMAG:%.*]] = load double, ptr [[DC1_IMAGP3]], align 8 +// HOST-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DP_ADDR]], align 8 +// HOST-NEXT: store double [[DC1_REAL]], ptr [[TMP2]], align 8 +// HOST-NEXT: [[FC_ADDR_REALP4:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[FC_ADDR]], i32 0, i32 0 +// HOST-NEXT: [[FC_ADDR_REAL5:%.*]] = load float, ptr [[FC_ADDR_REALP4]], align 4 +// HOST-NEXT: [[FC_ADDR_IMAGP6:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[FC_ADDR]], i32 0, i32 1 +// HOST-NEXT: [[FC_ADDR_IMAG7:%.*]] = load float, ptr [[FC_ADDR_IMAGP6]], align 4 +// HOST-NEXT: [[TMP3:%.*]] = load ptr, ptr [[FP_ADDR]], align 8 +// HOST-NEXT: store float [[FC_ADDR_IMAG7]], ptr [[TMP3]], align 4 +// HOST-NEXT: [[DC1_REALP8:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC1]], i32 0, i32 0 +// HOST-NEXT: [[DC1_REAL9:%.*]] = load double, ptr [[DC1_REALP8]], align 8 +// HOST-NEXT: [[DC1_IMAGP10:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC1]], i32 0, i32 1 +// HOST-NEXT: [[DC1_IMAG11:%.*]] = load double, ptr [[DC1_IMAGP10]], align 8 +// HOST-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DP_ADDR]], align 8 +// HOST-NEXT: store double [[DC1_IMAG11]], ptr [[TMP4]], align 8 +// HOST-NEXT: [[FC_ADDR_REALP12:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[FC_ADDR]], i32 0, i32 0 +// HOST-NEXT: [[FC_ADDR_REAL13:%.*]] = load float, ptr [[FC_ADDR_REALP12]], align 4 +// HOST-NEXT: [[FC_ADDR_IMAGP14:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[FC_ADDR]], i32 0, i32 1 +// HOST-NEXT: [[FC_ADDR_IMAG15:%.*]] = load float, ptr [[FC_ADDR_IMAGP14]], align 4 +// HOST-NEXT: [[NEG:%.*]] = fneg float [[FC_ADDR_IMAG15]] +// HOST-NEXT: [[TMP5:%.*]] = load ptr, ptr [[FCP_ADDR]], align 8 +// HOST-NEXT: [[DOTREALP16:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[TMP5]], i32 0, i32 0 +// HOST-NEXT: [[DOTIMAGP17:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[TMP5]], i32 0, i32 1 +// HOST-NEXT: store float [[FC_ADDR_REAL13]], ptr [[DOTREALP16]], align 4 +// HOST-NEXT: store float [[NEG]], ptr [[DOTIMAGP17]], align 4 +// HOST-NEXT: [[DC1_REALP18:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC1]], i32 0, i32 0 +// HOST-NEXT: [[DC1_REAL19:%.*]] = load double, ptr [[DC1_REALP18]], align 8 +// HOST-NEXT: [[DC1_IMAGP20:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC1]], i32 0, i32 1 +// HOST-NEXT: [[DC1_IMAG21:%.*]] = load double, ptr [[DC1_IMAGP20]], align 8 +// HOST-NEXT: [[NEG22:%.*]] = fneg double [[DC1_IMAG21]] +// HOST-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DCP_ADDR]], align 8 +// HOST-NEXT: [[DOTREALP23:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[TMP6]], i32 0, i32 0 +// HOST-NEXT: [[DOTIMAGP24:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[TMP6]], i32 0, i32 1 +// HOST-NEXT: store double [[DC1_REAL19]], ptr [[DOTREALP23]], align 8 +// HOST-NEXT: store double [[NEG22]], ptr [[DOTIMAGP24]], align 8 +// HOST-NEXT: ret void +// +// +// DEVICE-LABEL: define weak_odr protected amdgpu_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_test_complex_builtins_l39( +// DEVICE-SAME: ptr noundef [[FP:%.*]], i64 noundef [[FC:%.*]], ptr noundef [[DP:%.*]], ptr noundef nonnull align 8 dereferenceable(16) [[DC:%.*]], ptr noundef [[FCP:%.*]], ptr noundef [[DCP:%.*]], ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0:[0-9]+]] { +// DEVICE-NEXT: [[ENTRY:.*:]] +// DEVICE-NEXT: [[FP_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// DEVICE-NEXT: [[FC_ADDR:%.*]] = alloca i64, align 8, addrspace(5) +// DEVICE-NEXT: [[DP_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// DEVICE-NEXT: [[DC_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// DEVICE-NEXT: [[FCP_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// DEVICE-NEXT: [[DCP_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// DEVICE-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// DEVICE-NEXT: [[DC1:%.*]] = alloca { double, double }, align 8, addrspace(5) +// DEVICE-NEXT: [[FP_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[FP_ADDR]] to ptr +// DEVICE-NEXT: [[FC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[FC_ADDR]] to ptr +// DEVICE-NEXT: [[DP_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DP_ADDR]] to ptr +// DEVICE-NEXT: [[DC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DC_ADDR]] to ptr +// DEVICE-NEXT: [[FCP_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[FCP_ADDR]] to ptr +// DEVICE-NEXT: [[DCP_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DCP_ADDR]] to ptr +// DEVICE-NEXT: [[DYN_PTR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DYN_PTR_ADDR]] to ptr +// DEVICE-NEXT: [[DC1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DC1]] to ptr +// DEVICE-NEXT: store ptr [[FP]], ptr [[FP_ADDR_ASCAST]], align 8 +// DEVICE-NEXT: store i64 [[FC]], ptr [[FC_ADDR_ASCAST]], align 8 +// DEVICE-NEXT: store ptr [[DP]], ptr [[DP_ADDR_ASCAST]], align 8 +// DEVICE-NEXT: store ptr [[DC]], ptr [[DC_ADDR_ASCAST]], align 8 +// DEVICE-NEXT: store ptr [[FCP]], ptr [[FCP_ADDR_ASCAST]], align 8 +// DEVICE-NEXT: store ptr [[DCP]], ptr [[DCP_ADDR_ASCAST]], align 8 +// DEVICE-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR_ASCAST]], align 8 +// DEVICE-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DC_ADDR_ASCAST]], align 8, !nonnull [[META5:![0-9]+]], !align [[META6:![0-9]+]] +// DEVICE-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_test_complex_builtins_l39_kernel_environment to ptr), ptr [[DYN_PTR]]) +// DEVICE-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 +// DEVICE-NEXT: br i1 [[EXEC_USER_CODE]], label %[[USER_CODE_ENTRY:.*]], label %[[WORKER_EXIT:.*]] +// DEVICE: [[USER_CODE_ENTRY]]: +// DEVICE-NEXT: [[DOTREALP:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[TMP0]], i32 0, i32 0 +// DEVICE-NEXT: [[DOTREAL:%.*]] = load double, ptr [[DOTREALP]], align 8 +// DEVICE-NEXT: [[DOTIMAGP:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[TMP0]], i32 0, i32 1 +// DEVICE-NEXT: [[DOTIMAG:%.*]] = load double, ptr [[DOTIMAGP]], align 8 +// DEVICE-NEXT: [[DC1_ASCAST_REALP:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC1_ASCAST]], i32 0, i32 0 +// DEVICE-NEXT: [[DC1_ASCAST_IMAGP:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC1_ASCAST]], i32 0, i32 1 +// DEVICE-NEXT: store double [[DOTREAL]], ptr [[DC1_ASCAST_REALP]], align 8 +// DEVICE-NEXT: store double [[DOTIMAG]], ptr [[DC1_ASCAST_IMAGP]], align 8 +// DEVICE-NEXT: [[FC_ADDR_ASCAST_REALP:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[FC_ADDR_ASCAST]], i32 0, i32 0 +// DEVICE-NEXT: [[FC_ADDR_ASCAST_REAL:%.*]] = load float, ptr [[FC_ADDR_ASCAST_REALP]], align 4 +// DEVICE-NEXT: [[FC_ADDR_ASCAST_IMAGP:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[FC_ADDR_ASCAST]], i32 0, i32 1 +// DEVICE-NEXT: [[FC_ADDR_ASCAST_IMAG:%.*]] = load float, ptr [[FC_ADDR_ASCAST_IMAGP]], align 4 +// DEVICE-NEXT: [[TMP2:%.*]] = load ptr, ptr [[FP_ADDR_ASCAST]], align 8 +// DEVICE-NEXT: store float [[FC_ADDR_ASCAST_REAL]], ptr [[TMP2]], align 4 +// DEVICE-NEXT: [[DC1_ASCAST_REALP2:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC1_ASCAST]], i32 0, i32 0 +// DEVICE-NEXT: [[DC1_ASCAST_REAL:%.*]] = load double, ptr [[DC1_ASCAST_REALP2]], align 8 +// DEVICE-NEXT: [[DC1_ASCAST_IMAGP3:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC1_ASCAST]], i32 0, i32 1 +// DEVICE-NEXT: [[DC1_ASCAST_IMAG:%.*]] = load double, ptr [[DC1_ASCAST_IMAGP3]], align 8 +// DEVICE-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DP_ADDR_ASCAST]], align 8 +// DEVICE-NEXT: store double [[DC1_ASCAST_REAL]], ptr [[TMP3]], align 8 +// DEVICE-NEXT: [[FC_ADDR_ASCAST_REALP4:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[FC_ADDR_ASCAST]], i32 0, i32 0 +// DEVICE-NEXT: [[FC_ADDR_ASCAST_REAL5:%.*]] = load float, ptr [[FC_ADDR_ASCAST_REALP4]], align 4 +// DEVICE-NEXT: [[FC_ADDR_ASCAST_IMAGP6:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[FC_ADDR_ASCAST]], i32 0, i32 1 +// DEVICE-NEXT: [[FC_ADDR_ASCAST_IMAG7:%.*]] = load float, ptr [[FC_ADDR_ASCAST_IMAGP6]], align 4 +// DEVICE-NEXT: [[TMP4:%.*]] = load ptr, ptr [[FP_ADDR_ASCAST]], align 8 +// DEVICE-NEXT: store float [[FC_ADDR_ASCAST_IMAG7]], ptr [[TMP4]], align 4 +// DEVICE-NEXT: [[DC1_ASCAST_REALP8:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC1_ASCAST]], i32 0, i32 0 +// DEVICE-NEXT: [[DC1_ASCAST_REAL9:%.*]] = load double, ptr [[DC1_ASCAST_REALP8]], align 8 +// DEVICE-NEXT: [[DC1_ASCAST_IMAGP10:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC1_ASCAST]], i32 0, i32 1 +// DEVICE-NEXT: [[DC1_ASCAST_IMAG11:%.*]] = load double, ptr [[DC1_ASCAST_IMAGP10]], align 8 +// DEVICE-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DP_ADDR_ASCAST]], align 8 +// DEVICE-NEXT: store double [[DC1_ASCAST_IMAG11]], ptr [[TMP5]], align 8 +// DEVICE-NEXT: [[FC_ADDR_ASCAST_REALP12:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[FC_ADDR_ASCAST]], i32 0, i32 0 +// DEVICE-NEXT: [[FC_ADDR_ASCAST_REAL13:%.*]] = load float, ptr [[FC_ADDR_ASCAST_REALP12]], align 4 +// DEVICE-NEXT: [[FC_ADDR_ASCAST_IMAGP14:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[FC_ADDR_ASCAST]], i32 0, i32 1 +// DEVICE-NEXT: [[FC_ADDR_ASCAST_IMAG15:%.*]] = load float, ptr [[FC_ADDR_ASCAST_IMAGP14]], align 4 +// DEVICE-NEXT: [[NEG:%.*]] = fneg float [[FC_ADDR_ASCAST_IMAG15]] +// DEVICE-NEXT: [[TMP6:%.*]] = load ptr, ptr [[FCP_ADDR_ASCAST]], align 8 +// DEVICE-NEXT: [[DOTREALP16:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[TMP6]], i32 0, i32 0 +// DEVICE-NEXT: [[DOTIMAGP17:%.*]] = getelementptr inbounds nuw { float, float }, ptr [[TMP6]], i32 0, i32 1 +// DEVICE-NEXT: store float [[FC_ADDR_ASCAST_REAL13]], ptr [[DOTREALP16]], align 4 +// DEVICE-NEXT: store float [[NEG]], ptr [[DOTIMAGP17]], align 4 +// DEVICE-NEXT: [[DC1_ASCAST_REALP18:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC1_ASCAST]], i32 0, i32 0 +// DEVICE-NEXT: [[DC1_ASCAST_REAL19:%.*]] = load double, ptr [[DC1_ASCAST_REALP18]], align 8 +// DEVICE-NEXT: [[DC1_ASCAST_IMAGP20:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[DC1_ASCAST]], i32 0, i32 1 +// DEVICE-NEXT: [[DC1_ASCAST_IMAG21:%.*]] = load double, ptr [[DC1_ASCAST_IMAGP20]], align 8 +// DEVICE-NEXT: [[NEG22:%.*]] = fneg double [[DC1_ASCAST_IMAG21]] +// DEVICE-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DCP_ADDR_ASCAST]], align 8 +// DEVICE-NEXT: [[DOTREALP23:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[TMP7]], i32 0, i32 0 +// DEVICE-NEXT: [[DOTIMAGP24:%.*]] = getelementptr inbounds nuw { double, double }, ptr [[TMP7]], i32 0, i32 1 +// DEVICE-NEXT: store double [[DC1_ASCAST_REAL19]], ptr [[DOTREALP23]], align 8 +// DEVICE-NEXT: store double [[NEG22]], ptr [[DOTIMAGP24]], align 8 +// DEVICE-NEXT: call void @__kmpc_target_deinit() +// DEVICE-NEXT: ret void +// DEVICE: [[WORKER_EXIT]]: +// DEVICE-NEXT: ret void +// +//. +// HOST: [[META4]] = !{} +// HOST: [[META5]] = !{i64 8} +//. +// DEVICE: [[META5]] = !{} +// DEVICE: [[META6]] = !{i64 8} +//. _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
