Skip to content

Commit 958d599

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web'
2 parents 4e46bd3 + 44136bd commit 958d599

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

54 files changed

+2427
-351
lines changed

‎.github/workflows/clang-format.yml

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -14,11 +14,11 @@ jobs:
1414
fetch-depth: 2
1515

1616
- name: Get clang-format first
17-
run: sudo apt-get install -yqq clang-format-9
17+
run: sudo apt-get install -yqq clang-format-10
1818

1919
- name: Run clang-format for the patch
2020
run: |
21-
git diff -U0 --no-color ${GITHUB_SHA}^1 ${GITHUB_SHA} -- | ./clang/tools/clang-format/clang-format-diff.py -p1 -binary clang-format-9 > ./clang-format.patch
21+
git diff -U0 --no-color ${GITHUB_SHA}^1 ${GITHUB_SHA} -- | ./clang/tools/clang-format/clang-format-diff.py -p1 -binary clang-format-10 > ./clang-format.patch
2222
2323
# Add patch with formatting fixes to CI job artifacts
2424
- uses: actions/upload-artifact@v1

‎clang/lib/CodeGen/BackendUtil.cpp

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -41,6 +41,7 @@
4141
#include "llvm/Passes/PassBuilder.h"
4242
#include "llvm/Passes/PassPlugin.h"
4343
#include "llvm/Passes/StandardInstrumentations.h"
44+
#include "llvm/SYCLLowerIR/LowerESIMD.h"
4445
#include "llvm/Support/BuryPointer.h"
4546
#include "llvm/Support/CommandLine.h"
4647
#include "llvm/Support/MemoryBuffer.h"
@@ -786,6 +787,25 @@ void EmitAssemblyHelper::CreatePasses(legacy::PassManager &MPM,
786787

787788
PMBuilder.populateFunctionPassManager(FPM);
788789
PMBuilder.populateModulePassManager(MPM);
790+
791+
// Customize the tail of the module passes list for the ESIMD extension.
792+
if (LangOpts.SYCLIsDevice && LangOpts.SYCLExplicitSIMD &&
793+
CodeGenOpts.OptimizationLevel != 0) {
794+
MPM.add(createESIMDLowerVecArgPass());
795+
MPM.add(createESIMDLowerLoadStorePass());
796+
MPM.add(createSROAPass());
797+
MPM.add(createEarlyCSEPass(true));
798+
MPM.add(createInstructionCombiningPass());
799+
MPM.add(createDeadCodeEliminationPass());
800+
MPM.add(createFunctionInliningPass(
801+
CodeGenOpts.OptimizationLevel, CodeGenOpts.OptimizeSize,
802+
(!CodeGenOpts.SampleProfileFile.empty() &&
803+
CodeGenOpts.PrepareForThinLTO)));
804+
MPM.add(createSROAPass());
805+
MPM.add(createEarlyCSEPass(true));
806+
MPM.add(createInstructionCombiningPass());
807+
MPM.add(createDeadCodeEliminationPass());
808+
}
789809
}
790810

791811
static void setCommandLineOpts(const CodeGenOptions &CodeGenOpts) {
@@ -880,6 +900,11 @@ void EmitAssemblyHelper::EmitAssembly(BackendAction Action,
880900
PerFunctionPasses.add(
881901
createTargetTransformInfoWrapperPass(getTargetIRAnalysis()));
882902

903+
// ESIMD extension always requires lowering of certain IR constructs, such as
904+
// ESIMD C++ intrinsics, as the last FE step.
905+
if (LangOpts.SYCLIsDevice && LangOpts.SYCLExplicitSIMD)
906+
PerModulePasses.add(createSYCLLowerESIMDPass());
907+
883908
CreatePasses(PerModulePasses, PerFunctionPasses);
884909

885910
legacy::PassManager CodeGenPasses;

‎clang/lib/CodeGen/CodeGenModule.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -3887,13 +3887,13 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) {
38873887
}
38883888

38893889
if (LangOpts.SYCLIsDevice) {
3890+
if (!D)
3891+
return LangAS::opencl_global;
38903892
auto *Scope = D->getAttr<SYCLScopeAttr>();
3891-
38923893
if (Scope && Scope->isWorkGroup())
38933894
return LangAS::opencl_local;
3894-
if (!D || D->getType().getAddressSpace() == LangAS::Default) {
3895+
if (D->getType().getAddressSpace() == LangAS::Default)
38953896
return LangAS::opencl_global;
3896-
}
38973897
}
38983898

38993899
if (LangOpts.CUDA && LangOpts.CUDAIsDevice) {

‎clang/lib/Frontend/CompilerInvocation.cpp

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -816,10 +816,14 @@ static bool ParseCodeGenArgs(CodeGenOptions &Opts, ArgList &Args, InputKind IK,
816816
Args.getLastArg(OPT_emit_llvm_uselists, OPT_no_emit_llvm_uselists))
817817
Opts.EmitLLVMUseLists = A->getOption().getID() == OPT_emit_llvm_uselists;
818818

819+
// ESIMD GPU Back-end requires optimized IR
820+
bool IsSyclESIMD = Args.hasFlag(options::OPT_fsycl_esimd,
821+
options::OPT_fno_sycl_esimd, false);
822+
819823
Opts.DisableLLVMPasses =
820824
Args.hasArg(OPT_disable_llvm_passes) ||
821825
(Args.hasArg(OPT_fsycl_is_device) && Triple.isSPIR() &&
822-
!Args.hasArg(OPT_fsycl_enable_optimizations));
826+
!Args.hasArg(OPT_fsycl_enable_optimizations) && !IsSyclESIMD);
823827
Opts.DisableLifetimeMarkers = Args.hasArg(OPT_disable_lifetimemarkers);
824828

825829
const llvm::Triple::ArchType DebugEntryValueArchs[] = {

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 49 additions & 45 deletions
Original file line numberDiff line numberDiff line change
@@ -510,52 +510,25 @@ class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
510510

511511
if (auto *A = FD->getAttr<IntelReqdSubGroupSizeAttr>())
512512
Attrs.insert(A);
513+
513514
if (auto *A = FD->getAttr<ReqdWorkGroupSizeAttr>())
514515
Attrs.insert(A);
515-
// Allow the following kernel attributes only on lambda functions and
516-
// function objects that are called directly from a kernel (i.e. the one
517-
// passed to the parallel_for function). For all other cases,
518-
// emit a warning and ignore.
519-
if (auto *A = FD->getAttr<SYCLIntelKernelArgsRestrictAttr>()) {
520-
if (ParentFD == SYCLKernel) {
521-
Attrs.insert(A);
522-
} else {
523-
SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A;
524-
FD->dropAttr<SYCLIntelKernelArgsRestrictAttr>();
525-
}
526-
}
527-
if (auto *A = FD->getAttr<SYCLIntelNumSimdWorkItemsAttr>()) {
528-
if (ParentFD == SYCLKernel) {
529-
Attrs.insert(A);
530-
} else {
531-
SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A;
532-
FD->dropAttr<SYCLIntelNumSimdWorkItemsAttr>();
533-
}
534-
}
535-
if (auto *A = FD->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
536-
if (ParentFD == SYCLKernel) {
537-
Attrs.insert(A);
538-
} else {
539-
SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A;
540-
FD->dropAttr<SYCLIntelMaxWorkGroupSizeAttr>();
541-
}
542-
}
543-
if (auto *A = FD->getAttr<SYCLIntelMaxGlobalWorkDimAttr>()) {
544-
if (ParentFD == SYCLKernel) {
545-
Attrs.insert(A);
546-
} else {
547-
SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A;
548-
FD->dropAttr<SYCLIntelMaxGlobalWorkDimAttr>();
549-
}
550-
}
551-
if (auto *A = FD->getAttr<SYCLIntelNoGlobalWorkOffsetAttr>()) {
552-
if (ParentFD == SYCLKernel) {
553-
Attrs.insert(A);
554-
} else {
555-
SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A;
556-
FD->dropAttr<SYCLIntelNoGlobalWorkOffsetAttr>();
557-
}
558-
}
516+
517+
if (auto *A = FD->getAttr<SYCLIntelKernelArgsRestrictAttr>())
518+
Attrs.insert(A);
519+
520+
if (auto *A = FD->getAttr<SYCLIntelNumSimdWorkItemsAttr>())
521+
Attrs.insert(A);
522+
523+
if (auto *A = FD->getAttr<SYCLIntelMaxWorkGroupSizeAttr>())
524+
Attrs.insert(A);
525+
526+
if (auto *A = FD->getAttr<SYCLIntelMaxGlobalWorkDimAttr>())
527+
Attrs.insert(A);
528+
529+
if (auto *A = FD->getAttr<SYCLIntelNoGlobalWorkOffsetAttr>())
530+
Attrs.insert(A);
531+
559532
if (auto *A = FD->getAttr<SYCLSimdAttr>())
560533
Attrs.insert(A);
561534
// Propagate the explicit SIMD attribute through call graph - it is used
@@ -2051,6 +2024,38 @@ void Sema::MarkDevice(void) {
20512024
Diag(Attr->getLocation(), diag::note_conflicting_attribute);
20522025
SYCLKernel->setInvalidDecl();
20532026
}
2027+
} else if (auto *Existing =
2028+
SYCLKernel->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
2029+
if (Existing->getXDim() < Attr->getXDim() ||
2030+
Existing->getYDim() < Attr->getYDim() ||
2031+
Existing->getZDim() < Attr->getZDim()) {
2032+
Diag(SYCLKernel->getLocation(),
2033+
diag::err_conflicting_sycl_kernel_attributes);
2034+
Diag(Existing->getLocation(), diag::note_conflicting_attribute);
2035+
Diag(Attr->getLocation(), diag::note_conflicting_attribute);
2036+
SYCLKernel->setInvalidDecl();
2037+
} else {
2038+
SYCLKernel->addAttr(A);
2039+
}
2040+
} else {
2041+
SYCLKernel->addAttr(A);
2042+
}
2043+
break;
2044+
}
2045+
case attr::Kind::SYCLIntelMaxWorkGroupSize: {
2046+
auto *Attr = cast<SYCLIntelMaxWorkGroupSizeAttr>(A);
2047+
if (auto *Existing = SYCLKernel->getAttr<ReqdWorkGroupSizeAttr>()) {
2048+
if (Existing->getXDim() > Attr->getXDim() ||
2049+
Existing->getYDim() > Attr->getYDim() ||
2050+
Existing->getZDim() > Attr->getZDim()) {
2051+
Diag(SYCLKernel->getLocation(),
2052+
diag::err_conflicting_sycl_kernel_attributes);
2053+
Diag(Existing->getLocation(), diag::note_conflicting_attribute);
2054+
Diag(Attr->getLocation(), diag::note_conflicting_attribute);
2055+
SYCLKernel->setInvalidDecl();
2056+
} else {
2057+
SYCLKernel->addAttr(A);
2058+
}
20542059
} else {
20552060
SYCLKernel->addAttr(A);
20562061
}
@@ -2059,7 +2064,6 @@ void Sema::MarkDevice(void) {
20592064
case attr::Kind::SYCLIntelKernelArgsRestrict:
20602065
case attr::Kind::SYCLIntelNumSimdWorkItems:
20612066
case attr::Kind::SYCLIntelMaxGlobalWorkDim:
2062-
case attr::Kind::SYCLIntelMaxWorkGroupSize:
20632067
case attr::Kind::SYCLIntelNoGlobalWorkOffset:
20642068
case attr::Kind::SYCLSimd: {
20652069
if ((A->getKind() == attr::Kind::SYCLSimd) && KernelBody &&

‎clang/test/CodeGenSYCL/esimd-private-global.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,6 @@ __attribute__((opencl_private)) __attribute__((register_num(17))) int vc;
99

1010
SYCL_EXTERNAL void init_vc(int x) {
1111
vc = x;
12-
// CHECK: store i32 %0, i32* @vc
12+
// CHECK: store i32 %{{[0-9a-zA-Z_]+}}, i32* @vc
1313
}
1414
// CHECK: attributes #0 = {{.*"VCByteOffset"="17".*"VCVolatile"}}

‎clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp

Lines changed: 54 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,54 @@
1+
// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -verify
2+
// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -DTRIGGER_ERROR -verify
3+
// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 | FileCheck %s
4+
5+
#ifndef TRIGGER_ERROR
6+
[[intelfpga::no_global_work_offset]] void not_direct_one() {} // expected-no-diagnostics
7+
8+
[[intel::reqd_sub_group_size(1)]] void func_one() {
9+
not_direct_one();
10+
}
11+
12+
#else
13+
[[cl::reqd_work_group_size(2, 2, 2)]] void not_direct_two() {} // expected-note {{conflicting attribute is here}}
14+
15+
[[intelfpga::max_work_group_size(1, 1, 1)]] // expected-note {{conflicting attribute is here}}
16+
void
17+
func_two() {
18+
not_direct_two();
19+
}
20+
21+
[[cl::reqd_work_group_size(4, 4, 4)]] // expected-note 2 {{conflicting attribute is here}}
22+
void
23+
func_three() {
24+
not_direct_two();
25+
}
26+
#endif
27+
28+
template <typename Name, typename Type>
29+
[[clang::sycl_kernel]] void __my_kernel__(Type bar) {
30+
bar();
31+
#ifndef TRIGGER_ERROR
32+
func_one();
33+
#else
34+
func_two();
35+
func_three();
36+
#endif
37+
}
38+
39+
template <typename Name, typename Type>
40+
void parallel_for(Type lambda) {
41+
__my_kernel__<Name>(lambda);
42+
}
43+
44+
void invoke_foo2() {
45+
#ifndef TRIGGER_ERROR
46+
// CHECK-LABEL: FunctionDecl {{.*}} invoke_foo2 'void ()'
47+
// CHECK: `-FunctionDecl {{.*}}KernelName 'void ()'
48+
// CHECK: -IntelReqdSubGroupSizeAttr {{.*}}
49+
// CHECK: `-SYCLIntelNoGlobalWorkOffsetAttr {{.*}} Enabled
50+
parallel_for<class KernelName>([]() {});
51+
#else
52+
parallel_for<class KernelName>([]() {}); // expected-error 2 {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}}
53+
#endif
54+
}

‎clang/test/SemaSYCL/intel-max-global-work-dim.cpp

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -20,8 +20,7 @@ void foo() {
2020

2121
#else // __SYCL_DEVICE_ONLY__
2222

23-
[[intelfpga::max_global_work_dim(2)]] // expected-warning{{'max_global_work_dim' attribute ignored}}
24-
void func_ignore() {}
23+
[[intelfpga::max_global_work_dim(2)]] void func_do_not_ignore() {}
2524

2625
struct FuncObj {
2726
[[intelfpga::max_global_work_dim(1)]]
@@ -68,9 +67,9 @@ int main() {
6867
[]() [[intelfpga::max_global_work_dim(2)]] {});
6968

7069
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
71-
// CHECK-NOT: SYCLIntelMaxGlobalWorkDimAttr {{.*}}
70+
// CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}}
7271
kernel<class test_kernel3>(
73-
[]() {func_ignore();});
72+
[]() { func_do_not_ignore(); });
7473

7574
kernel<class test_kernel4>(
7675
TRIFuncObjGood1());

‎clang/test/SemaSYCL/intel-max-work-group-size.cpp

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -20,8 +20,7 @@ void foo() {
2020

2121
#else // __SYCL_DEVICE_ONLY__
2222

23-
[[intelfpga::max_work_group_size(2, 2, 2)]] // expected-warning{{'max_work_group_size' attribute ignored}}
24-
void func_ignore() {}
23+
[[intelfpga::max_work_group_size(2, 2, 2)]] void func_do_not_ignore() {}
2524

2625
struct FuncObj {
2726
[[intelfpga::max_work_group_size(4, 4, 4)]]
@@ -53,9 +52,9 @@ int main() {
5352
[]() [[intelfpga::max_work_group_size(8, 8, 8)]] {});
5453

5554
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
56-
// CHECK-NOT: SYCLIntelMaxWorkGroupSizeAttr {{.*}}
55+
// CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}}
5756
kernel<class test_kernel3>(
58-
[]() {func_ignore();});
57+
[]() { func_do_not_ignore(); });
5958

6059
#ifdef TRIGGER_ERROR
6160
[[intelfpga::max_work_group_size(1, 1, 1)]] int Var = 0; // expected-error{{'max_work_group_size' attribute only applies to functions}}

‎clang/test/SemaSYCL/intel-restrict.cpp

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,7 @@
11
// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -DCHECKDIAG -verify
22
// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 | FileCheck %s
33

4-
[[intel::kernel_args_restrict]] // expected-warning{{'kernel_args_restrict' attribute ignored}}
5-
void func_ignore() {}
4+
[[intel::kernel_args_restrict]] void func_do_not_ignore() {}
65

76
struct FuncObj {
87
[[intel::kernel_args_restrict]]
@@ -29,7 +28,7 @@ int main() {
2928
[]() [[intel::kernel_args_restrict]] {});
3029

3130
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
32-
// CHECK-NOT: SYCLIntelKernelArgsRestrictAttr
31+
// CHECK: SYCLIntelKernelArgsRestrictAttr
3332
kernel<class test_kernel3>(
34-
[]() {func_ignore();});
33+
[]() { func_do_not_ignore(); });
3534
}

‎clang/test/SemaSYCL/num_simd_work_items.cpp

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -20,8 +20,7 @@ void foo() {
2020

2121
#else // __SYCL_DEVICE_ONLY__
2222

23-
[[intelfpga::num_simd_work_items(2)]] // expected-warning{{'num_simd_work_items' attribute ignored}}
24-
void func_ignore() {}
23+
[[intelfpga::num_simd_work_items(2)]] void func_do_not_ignore() {}
2524

2625
struct FuncObj {
2726
[[intelfpga::num_simd_work_items(42)]]
@@ -45,9 +44,9 @@ int main() {
4544
[]() [[intelfpga::num_simd_work_items(8)]] {});
4645

4746
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
48-
// CHECK-NOT: SYCLIntelNumSimdWorkItemsAttr {{.*}} 2
47+
// CHECK: SYCLIntelNumSimdWorkItemsAttr {{.*}} 2
4948
kernel<class test_kernel3>(
50-
[]() {func_ignore();});
49+
[]() { func_do_not_ignore(); });
5150

5251
#ifdef TRIGGER_ERROR
5352
[[intelfpga::num_simd_work_items(0)]] int Var = 0; // expected-error{{'num_simd_work_items' attribute only applies to functions}}

‎clang/test/utils/update_cc_test_checks/Inputs/check-attributes.cpp.funcattrs.expected

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -17,12 +17,12 @@ struct ST {
1717
// CHECK-NEXT: [[S_ADDR:%.*]] = alloca %struct.ST*, align 8
1818
// CHECK-NEXT: store %struct.ST* [[S:%.*]], %struct.ST** [[S_ADDR]], align 8
1919
// CHECK-NEXT: [[TMP0:%.*]] = load %struct.ST*, %struct.ST** [[S_ADDR]], align 8
20-
// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [[STRUCT_ST:%.*]], %struct.ST* [[TMP0]], i64 1
21-
// CHECK-NEXT: [[Z:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ARRAYIDX]], i32 0, i32 2
20+
// CHECK-NEXT: [[PTRIDX:%.*]] = getelementptr inbounds [[STRUCT_ST:%.*]], %struct.ST* [[TMP0]], i64 1
21+
// CHECK-NEXT: [[Z:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[PTRIDX]], i32 0, i32 2
2222
// CHECK-NEXT: [[B:%.*]] = getelementptr inbounds [[STRUCT_RT:%.*]], %struct.RT* [[Z]], i32 0, i32 1
23-
// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [10 x [20 x i32]], [10 x [20 x i32]]* [[B]], i64 0, i64 5
24-
// CHECK-NEXT: [[ARRAYIDX2:%.*]] = getelementptr inbounds [20 x i32], [20 x i32]* [[ARRAYIDX1]], i64 0, i64 13
25-
// CHECK-NEXT: ret i32* [[ARRAYIDX2]]
23+
// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [20 x i32]], [10 x [20 x i32]]* [[B]], i64 0, i64 5
24+
// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [20 x i32], [20 x i32]* [[ARRAYIDX]], i64 0, i64 13
25+
// CHECK-NEXT: ret i32* [[ARRAYIDX1]]
2626
//
2727
int *foo(struct ST *s) {
2828
return &s[1].Z.B[5][13];

‎clang/test/utils/update_cc_test_checks/Inputs/check-attributes.cpp.plain.expected

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -16,12 +16,12 @@ struct ST {
1616
// CHECK-NEXT: [[S_ADDR:%.*]] = alloca %struct.ST*, align 8
1717
// CHECK-NEXT: store %struct.ST* [[S:%.*]], %struct.ST** [[S_ADDR]], align 8
1818
// CHECK-NEXT: [[TMP0:%.*]] = load %struct.ST*, %struct.ST** [[S_ADDR]], align 8
19-
// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [[STRUCT_ST:%.*]], %struct.ST* [[TMP0]], i64 1
20-
// CHECK-NEXT: [[Z:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ARRAYIDX]], i32 0, i32 2
19+
// CHECK-NEXT: [[PTRIDX:%.*]] = getelementptr inbounds [[STRUCT_ST:%.*]], %struct.ST* [[TMP0]], i64 1
20+
// CHECK-NEXT: [[Z:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[PTRIDX]], i32 0, i32 2
2121
// CHECK-NEXT: [[B:%.*]] = getelementptr inbounds [[STRUCT_RT:%.*]], %struct.RT* [[Z]], i32 0, i32 1
22-
// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [10 x [20 x i32]], [10 x [20 x i32]]* [[B]], i64 0, i64 5
23-
// CHECK-NEXT: [[ARRAYIDX2:%.*]] = getelementptr inbounds [20 x i32], [20 x i32]* [[ARRAYIDX1]], i64 0, i64 13
24-
// CHECK-NEXT: ret i32* [[ARRAYIDX2]]
22+
// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [20 x i32]], [10 x [20 x i32]]* [[B]], i64 0, i64 5
23+
// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [20 x i32], [20 x i32]* [[ARRAYIDX]], i64 0, i64 13
24+
// CHECK-NEXT: ret i32* [[ARRAYIDX1]]
2525
//
2626
int *foo(struct ST *s) {
2727
return &s[1].Z.B[5][13];

0 commit comments

Comments
 (0)