Skip to content

Commit 2c13dec

Browse files
AlexVlxAlexVlx
authored
[clang][llvm][SPIR-V] Explicitly encode native integer widths for SPIR-V (#110695)
SPIR-V doesn't currently encode "native" integer bit-widths in its datalayout(s). This is problematic as it leads to optimisation passes, such as InstCombine, getting ideas and e.g. shrinking to non byte-multiple integer types, which is not desirable and can lead to breakage further down in the toolchain. This patch addresses that by encoding `i8`, `i16`, `i32` and `i64` as native types for vanilla SPIR-V (the spec natively supports them), and `i32` and `i64` for AMDGCNSPIRV (where the hardware targets are known). We also set the stack alignment on the latter, as it is overaligned (32-bit vs 8-bit).
1 parent c50bb99 commit 2c13dec

File tree

5 files changed

+69
-34
lines changed

5 files changed

+69
-34
lines changed

‎clang/lib/Basic/Targets/SPIR.h

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -310,8 +310,8 @@ class LLVM_LIBRARY_VISIBILITY SPIRVTargetInfo : public BaseSPIRVTargetInfo {
310310

311311
// SPIR-V IDs are represented with a single 32-bit word.
312312
SizeType = TargetInfo::UnsignedInt;
313-
resetDataLayout("e-i64:64-v16:16-v24:32-v32:32-v48:64-"
314-
"v96:128-v192:256-v256:256-v512:512-v1024:1024-G1");
313+
resetDataLayout("e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-"
314+
"v256:256-v512:512-v1024:1024-n8:16:32:64-G1");
315315
}
316316

317317
void getTargetDefines(const LangOptions &Opts,
@@ -334,8 +334,8 @@ class LLVM_LIBRARY_VISIBILITY SPIRV32TargetInfo : public BaseSPIRVTargetInfo {
334334
// SPIR-V has core support for atomic ops, and Int32 is always available;
335335
// we take the maximum because it's possible the Host supports wider types.
336336
MaxAtomicInlineWidth = std::max<unsigned char>(MaxAtomicInlineWidth, 32);
337-
resetDataLayout("e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-"
338-
"v96:128-v192:256-v256:256-v512:512-v1024:1024-G1");
337+
resetDataLayout("e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-"
338+
"v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1");
339339
}
340340

341341
void getTargetDefines(const LangOptions &Opts,
@@ -358,8 +358,8 @@ class LLVM_LIBRARY_VISIBILITY SPIRV64TargetInfo : public BaseSPIRVTargetInfo {
358358
// SPIR-V has core support for atomic ops, and Int64 is always available;
359359
// we take the maximum because it's possible the Host supports wider types.
360360
MaxAtomicInlineWidth = std::max<unsigned char>(MaxAtomicInlineWidth, 64);
361-
resetDataLayout("e-i64:64-v16:16-v24:32-v32:32-v48:64-"
362-
"v96:128-v192:256-v256:256-v512:512-v1024:1024-G1");
361+
resetDataLayout("e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-"
362+
"v256:256-v512:512-v1024:1024-n8:16:32:64-G1");
363363
}
364364

365365
void getTargetDefines(const LangOptions &Opts,
@@ -384,8 +384,8 @@ class LLVM_LIBRARY_VISIBILITY SPIRV64AMDGCNTargetInfo final
384384
PtrDiffType = IntPtrType = TargetInfo::SignedLong;
385385
AddrSpaceMap = &SPIRDefIsGenMap;
386386

387-
resetDataLayout("e-i64:64-v16:16-v24:32-v32:32-v48:64-"
388-
"v96:128-v192:256-v256:256-v512:512-v1024:1024-G1-P4-A0");
387+
resetDataLayout("e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-"
388+
"v256:256-v512:512-v1024:1024-n32:64-S32-G1-P4-A0");
389389

390390
BFloat16Width = BFloat16Align = 16;
391391
BFloat16Format = &llvm::APFloat::BFloat();

‎clang/test/CodeGen/target-data.c

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -271,4 +271,4 @@
271271

272272
// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -o - -emit-llvm %s | \
273273
// RUN: FileCheck %s -check-prefix=AMDGPUSPIRV64
274-
// AMDGPUSPIRV64: target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-G1-P4-A0"
274+
// AMDGPUSPIRV64: target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n32:64-S32-G1-P4-A0"

‎clang/test/CodeGenOpenCL/builtins-amdgcn.cl

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -638,7 +638,7 @@ void test_get_workgroup_size(int d, global int *out)
638638

639639
// CHECK-LABEL: @test_get_grid_size(
640640
// CHECK: {{.*}}call align 4 dereferenceable(64){{.*}} ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
641-
// CHECK: getelementptr inbounds i8, ptr addrspace(4) %{{.*}}, i64 %.sink
641+
// CHECK: getelementptr inbounds i8, ptr addrspace(4) %{{.*}}, i64 %{{.+}}
642642
// CHECK: load i32, ptr addrspace(4) %{{.*}}, align 4, !invariant.load
643643
void test_get_grid_size(int d, global int *out)
644644
{

‎llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -55,14 +55,14 @@ static std::string computeDataLayout(const Triple &TT) {
5555
// memory model used for graphics: PhysicalStorageBuffer64. But it shouldn't
5656
// mean anything.
5757
if (Arch == Triple::spirv32)
58-
return "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-"
59-
"v96:128-v192:256-v256:256-v512:512-v1024:1024-G1";
58+
return "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-"
59+
"v256:256-v512:512-v1024:1024-n8:16:32:64-G1";
6060
if (TT.getVendor() == Triple::VendorType::AMD &&
6161
TT.getOS() == Triple::OSType::AMDHSA)
62-
return "e-i64:64-v16:16-v24:32-v32:32-v48:64-"
63-
"v96:128-v192:256-v256:256-v512:512-v1024:1024-G1-P4-A0";
64-
return "e-i64:64-v16:16-v24:32-v32:32-v48:64-"
65-
"v96:128-v192:256-v256:256-v512:512-v1024:1024-G1";
62+
return "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-"
63+
"v512:512-v1024:1024-n32:64-S32-G1-P4-A0";
64+
return "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-"
65+
"v512:512-v1024:1024-n8:16:32:64-G1";
6666
}
6767

6868
static Reloc::Model getEffectiveRelocModel(std::optional<Reloc::Model> RM) {

‎llvm/test/CodeGen/SPIRV/optimizations/add-check-overflow.ll

Lines changed: 53 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -1,41 +1,76 @@
11
; This test aims to check ability to support "Arithmetic with Overflow" intrinsics
22
; in the special case when those intrinsics are being generated by the CodeGenPrepare;
3-
; pass during translations with optimization (note -O3 in llc arguments).
3+
; pass during translations with optimization (note -disable-lsr, to inhibit
4+
; strength reduction pre-empting with a more preferable match for this pattern
5+
; in llc arguments).
46

57
; RUN: llc -O3 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
68
; RUN: %if spirv-tools %{ llc -O3 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
79

810
; RUN: llc -O3 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
911
; RUN: %if spirv-tools %{ llc -O3 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
1012

11-
; CHECK-DAG: OpName %[[Val:.*]] "math"
12-
; CHECK-DAG: OpName %[[IsOver:.*]] "ov"
13+
; RUN: llc -O3 -disable-lsr -mtriple=spirv32-unknown-unknown %s -o - | FileCheck --check-prefix=NOLSR %s
14+
; RUN: %if spirv-tools %{ llc -O3 -disable-lsr -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
15+
16+
; RUN: llc -O3 -disable-lsr -mtriple=spirv64-unknown-unknown %s -o - | FileCheck --check-prefix=NOLSR %s
17+
; RUN: %if spirv-tools %{ llc -O3 -disable-lsr -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
18+
19+
; CHECK-DAG: OpName %[[PhiRes:.*]] "lsr.iv"
20+
; CHECK-DAG: OpName %[[IsOver:.*]] "fl"
21+
; CHECK-DAG: OpName %[[Val:.*]] "lsr.iv.next"
1322
; CHECK-DAG: %[[Int:.*]] = OpTypeInt 32 0
1423
; CHECK-DAG: %[[Char:.*]] = OpTypeInt 8 0
1524
; CHECK-DAG: %[[PtrChar:.*]] = OpTypePointer Generic %[[Char]]
1625
; CHECK-DAG: %[[Bool:.*]] = OpTypeBool
17-
; CHECK-DAG: %[[Struct:.*]] = OpTypeStruct %[[Int]] %[[Int]]
1826
; CHECK-DAG: %[[Const1:.*]] = OpConstant %[[Int]] 1
27+
; CHECK-DAG: %[[Zero:.*]] = OpConstant %[[Int]] 0
1928
; CHECK-DAG: %[[Const42:.*]] = OpConstant %[[Char]] 42
20-
; CHECK-DAG: %[[Zero:.*]] = OpConstantNull %[[Int]]
2129

2230
; CHECK: OpFunction
2331
; CHECK: %[[A:.*]] = OpFunctionParameter %[[Int]]
2432
; CHECK: %[[Ptr:.*]] = OpFunctionParameter %[[PtrChar]]
25-
; CHECK: %[[#]] = OpLabel
26-
; CHECK: OpBranch %[[#]]
27-
; CHECK: %[[#]] = OpLabel
28-
; CHECK: %[[PhiRes:.*]] = OpPhi %[[Int]] %[[A]] %[[#]] %[[Val]] %[[#]]
29-
; CHECK: %[[AggRes:.*]] = OpIAddCarry %[[Struct]] %[[PhiRes]] %[[Const1]]
30-
; CHECK: %[[Val]] = OpCompositeExtract %[[Int]] %[[AggRes]] 0
31-
; CHECK: %[[Over:.*]] = OpCompositeExtract %[[Int]] %[[AggRes]] 1
32-
; CHECK: %[[IsOver]] = OpINotEqual %[[Bool:.*]] %[[Over]] %[[Zero]]
33-
; CHECK: OpBranchConditional %[[IsOver]] %[[#]] %[[#]]
34-
; CHECK: OpStore %[[Ptr]] %[[Const42]] Aligned 1
33+
; CHECK: %[[APlusOne:.*]] = OpIAdd %[[Int]] %[[A]] %[[Const1]]
34+
; CHECK: OpBranch %[[#]]
35+
; CHECK: [[#]] = OpLabel
36+
; CHECK: %[[PhiRes]] = OpPhi %[[Int]] %[[Val]] %[[#]] %[[APlusOne]] %[[#]]
37+
; CHECK: %[[IsOver]] = OpIEqual %[[Bool]] %[[#]] %[[#]]
38+
; CHECK: OpBranchConditional %[[IsOver]] %[[#]] %[[#]]
39+
; CHECK: [[#]] = OpLabel
40+
; CHECK: OpStore %[[Ptr]] %[[Const42]] Aligned 1
41+
; CHECK: [[Val]] = OpIAdd %[[Int]] %[[PhiRes]] %[[Const1]]
3542
; CHECK: OpBranch %[[#]]
36-
; CHECK: %[[#]] = OpLabel
37-
; CHECK: OpReturnValue %[[Val]]
38-
; CHECK: OpFunctionEnd
43+
; CHECK: [[#]] = OpLabel
44+
; OpReturnValue %[[PhiRes]]
45+
46+
; NOLSR-DAG: OpName %[[Val:.*]] "math"
47+
; NOLSR-DAG: OpName %[[IsOver:.*]] "ov"
48+
; NOLSR-DAG: %[[Int:.*]] = OpTypeInt 32 0
49+
; NOLSR-DAG: %[[Char:.*]] = OpTypeInt 8 0
50+
; NOLSR-DAG: %[[PtrChar:.*]] = OpTypePointer Generic %[[Char]]
51+
; NOLSR-DAG: %[[Bool:.*]] = OpTypeBool
52+
; NOLSR-DAG: %[[Struct:.*]] = OpTypeStruct %[[Int]] %[[Int]]
53+
; NOLSR-DAG: %[[Const1:.*]] = OpConstant %[[Int]] 1
54+
; NOLSR-DAG: %[[Const42:.*]] = OpConstant %[[Char]] 42
55+
; NOLSR-DAG: %[[Zero:.*]] = OpConstantNull %[[Int]]
56+
57+
; NOLSR: OpFunction
58+
; NOLSR: %[[A:.*]] = OpFunctionParameter %[[Int]]
59+
; NOLSR: %[[Ptr:.*]] = OpFunctionParameter %[[PtrChar]]
60+
; NOLSR: %[[#]] = OpLabel
61+
; NOLSR: OpBranch %[[#]]
62+
; NOLSR: %[[#]] = OpLabel
63+
; NOLSR: %[[PhiRes:.*]] = OpPhi %[[Int]] %[[A]] %[[#]] %[[Val]] %[[#]]
64+
; NOLSR: %[[AggRes:.*]] = OpIAddCarry %[[Struct]] %[[PhiRes]] %[[Const1]]
65+
; NOLSR: %[[Val]] = OpCompositeExtract %[[Int]] %[[AggRes]] 0
66+
; NOLSR: %[[Over:.*]] = OpCompositeExtract %[[Int]] %[[AggRes]] 1
67+
; NOLSR: %[[IsOver]] = OpINotEqual %[[Bool:.*]] %[[Over]] %[[Zero]]
68+
; NOLSR: OpBranchConditional %[[IsOver]] %[[#]] %[[#]]
69+
; NOLSR: OpStore %[[Ptr]] %[[Const42]] Aligned 1
70+
; NOLSR: OpBranch %[[#]]
71+
; NOLSR: %[[#]] = OpLabel
72+
; NOLSR: OpReturnValue %[[Val]]
73+
; NOLSR: OpFunctionEnd
3974

4075
define spir_func i32 @foo(i32 %a, ptr addrspace(4) %p) {
4176
entry:

0 commit comments

Comments
 (0)