[clang] 2c13dec - [clang][llvm][SPIR-V] Explicitly encode native integer widths for SPIR-V (#110695)
via cfe-commits
cfe-commits at lists.llvm.org
Tue Nov 5 07:26:12 PST 2024
Author: Alex Voicu
Date: 2024-11-05T17:26:08+02:00
New Revision: 2c13dec3284d019fdedf7913083ce82aa5cb97aa
URL: https://github.com/llvm/llvm-project/commit/2c13dec3284d019fdedf7913083ce82aa5cb97aa
DIFF: https://github.com/llvm/llvm-project/commit/2c13dec3284d019fdedf7913083ce82aa5cb97aa.diff
LOG: [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).
Added:
Modified:
clang/lib/Basic/Targets/SPIR.h
clang/test/CodeGen/target-data.c
clang/test/CodeGenOpenCL/builtins-amdgcn.cl
llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp
llvm/test/CodeGen/SPIRV/optimizations/add-check-overflow.ll
Removed:
################################################################################
diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h
index 30c7ac8d9c037f..85e4bd920d8535 100644
--- a/clang/lib/Basic/Targets/SPIR.h
+++ b/clang/lib/Basic/Targets/SPIR.h
@@ -310,8 +310,8 @@ class LLVM_LIBRARY_VISIBILITY SPIRVTargetInfo : public BaseSPIRVTargetInfo {
// SPIR-V IDs are represented with a single 32-bit word.
SizeType = TargetInfo::UnsignedInt;
- resetDataLayout("e-i64:64-v16:16-v24:32-v32:32-v48:64-"
- "v96:128-v192:256-v256:256-v512:512-v1024:1024-G1");
+ resetDataLayout("e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-"
+ "v256:256-v512:512-v1024:1024-n8:16:32:64-G1");
}
void getTargetDefines(const LangOptions &Opts,
@@ -334,8 +334,8 @@ class LLVM_LIBRARY_VISIBILITY SPIRV32TargetInfo : public BaseSPIRVTargetInfo {
// SPIR-V has core support for atomic ops, and Int32 is always available;
// we take the maximum because it's possible the Host supports wider types.
MaxAtomicInlineWidth = std::max<unsigned char>(MaxAtomicInlineWidth, 32);
- resetDataLayout("e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-"
- "v96:128-v192:256-v256:256-v512:512-v1024:1024-G1");
+ resetDataLayout("e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-"
+ "v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1");
}
void getTargetDefines(const LangOptions &Opts,
@@ -358,8 +358,8 @@ class LLVM_LIBRARY_VISIBILITY SPIRV64TargetInfo : public BaseSPIRVTargetInfo {
// SPIR-V has core support for atomic ops, and Int64 is always available;
// we take the maximum because it's possible the Host supports wider types.
MaxAtomicInlineWidth = std::max<unsigned char>(MaxAtomicInlineWidth, 64);
- resetDataLayout("e-i64:64-v16:16-v24:32-v32:32-v48:64-"
- "v96:128-v192:256-v256:256-v512:512-v1024:1024-G1");
+ resetDataLayout("e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-"
+ "v256:256-v512:512-v1024:1024-n8:16:32:64-G1");
}
void getTargetDefines(const LangOptions &Opts,
@@ -384,8 +384,8 @@ class LLVM_LIBRARY_VISIBILITY SPIRV64AMDGCNTargetInfo final
PtrDiffType = IntPtrType = TargetInfo::SignedLong;
AddrSpaceMap = &SPIRDefIsGenMap;
- resetDataLayout("e-i64:64-v16:16-v24:32-v32:32-v48:64-"
- "v96:128-v192:256-v256:256-v512:512-v1024:1024-G1-P4-A0");
+ resetDataLayout("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");
BFloat16Width = BFloat16Align = 16;
BFloat16Format = &llvm::APFloat::BFloat();
diff --git a/clang/test/CodeGen/target-data.c b/clang/test/CodeGen/target-data.c
index 26a1bf2a1a5740..2dc6ead54f5f3e 100644
--- a/clang/test/CodeGen/target-data.c
+++ b/clang/test/CodeGen/target-data.c
@@ -271,4 +271,4 @@
// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -o - -emit-llvm %s | \
// RUN: FileCheck %s -check-prefix=AMDGPUSPIRV64
-// 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"
+// 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"
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index bf5f2971cf118c..9132cc8a717e0f 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -638,7 +638,7 @@ void test_get_workgroup_size(int d, global int *out)
// CHECK-LABEL: @test_get_grid_size(
// CHECK: {{.*}}call align 4 dereferenceable(64){{.*}} ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
-// CHECK: getelementptr inbounds i8, ptr addrspace(4) %{{.*}}, i64 %.sink
+// CHECK: getelementptr inbounds i8, ptr addrspace(4) %{{.*}}, i64 %{{.+}}
// CHECK: load i32, ptr addrspace(4) %{{.*}}, align 4, !invariant.load
void test_get_grid_size(int d, global int *out)
{
diff --git a/llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp b/llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp
index 34854f31b3e387..194ce7c10bfd3f 100644
--- a/llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp
@@ -55,14 +55,14 @@ static std::string computeDataLayout(const Triple &TT) {
// memory model used for graphics: PhysicalStorageBuffer64. But it shouldn't
// mean anything.
if (Arch == Triple::spirv32)
- return "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-"
- "v96:128-v192:256-v256:256-v512:512-v1024:1024-G1";
+ return "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-"
+ "v256:256-v512:512-v1024:1024-n8:16:32:64-G1";
if (TT.getVendor() == Triple::VendorType::AMD &&
TT.getOS() == Triple::OSType::AMDHSA)
- return "e-i64:64-v16:16-v24:32-v32:32-v48:64-"
- "v96:128-v192:256-v256:256-v512:512-v1024:1024-G1-P4-A0";
- return "e-i64:64-v16:16-v24:32-v32:32-v48:64-"
- "v96:128-v192:256-v256:256-v512:512-v1024:1024-G1";
+ return "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";
+ return "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-"
+ "v512:512-v1024:1024-n8:16:32:64-G1";
}
static Reloc::Model getEffectiveRelocModel(std::optional<Reloc::Model> RM) {
diff --git a/llvm/test/CodeGen/SPIRV/optimizations/add-check-overflow.ll b/llvm/test/CodeGen/SPIRV/optimizations/add-check-overflow.ll
index 1a630f77a44c5d..e04678f802d7c7 100644
--- a/llvm/test/CodeGen/SPIRV/optimizations/add-check-overflow.ll
+++ b/llvm/test/CodeGen/SPIRV/optimizations/add-check-overflow.ll
@@ -1,6 +1,8 @@
; This test aims to check ability to support "Arithmetic with Overflow" intrinsics
; in the special case when those intrinsics are being generated by the CodeGenPrepare;
-; pass during translations with optimization (note -O3 in llc arguments).
+; pass during translations with optimization (note -disable-lsr, to inhibit
+; strength reduction pre-empting with a more preferable match for this pattern
+; in llc arguments).
; RUN: llc -O3 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
; RUN: %if spirv-tools %{ llc -O3 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
@@ -8,34 +10,67 @@
; RUN: llc -O3 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
; RUN: %if spirv-tools %{ llc -O3 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
-; CHECK-DAG: OpName %[[Val:.*]] "math"
-; CHECK-DAG: OpName %[[IsOver:.*]] "ov"
+; RUN: llc -O3 -disable-lsr -mtriple=spirv32-unknown-unknown %s -o - | FileCheck --check-prefix=NOLSR %s
+; RUN: %if spirv-tools %{ llc -O3 -disable-lsr -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; RUN: llc -O3 -disable-lsr -mtriple=spirv64-unknown-unknown %s -o - | FileCheck --check-prefix=NOLSR %s
+; RUN: %if spirv-tools %{ llc -O3 -disable-lsr -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; CHECK-DAG: OpName %[[PhiRes:.*]] "lsr.iv"
+; CHECK-DAG: OpName %[[IsOver:.*]] "fl"
+; CHECK-DAG: OpName %[[Val:.*]] "lsr.iv.next"
; CHECK-DAG: %[[Int:.*]] = OpTypeInt 32 0
; CHECK-DAG: %[[Char:.*]] = OpTypeInt 8 0
; CHECK-DAG: %[[PtrChar:.*]] = OpTypePointer Generic %[[Char]]
; CHECK-DAG: %[[Bool:.*]] = OpTypeBool
-; CHECK-DAG: %[[Struct:.*]] = OpTypeStruct %[[Int]] %[[Int]]
; CHECK-DAG: %[[Const1:.*]] = OpConstant %[[Int]] 1
+; CHECK-DAG: %[[Zero:.*]] = OpConstant %[[Int]] 0
; CHECK-DAG: %[[Const42:.*]] = OpConstant %[[Char]] 42
-; CHECK-DAG: %[[Zero:.*]] = OpConstantNull %[[Int]]
; CHECK: OpFunction
; CHECK: %[[A:.*]] = OpFunctionParameter %[[Int]]
; CHECK: %[[Ptr:.*]] = OpFunctionParameter %[[PtrChar]]
-; CHECK: %[[#]] = OpLabel
-; CHECK: OpBranch %[[#]]
-; CHECK: %[[#]] = OpLabel
-; CHECK: %[[PhiRes:.*]] = OpPhi %[[Int]] %[[A]] %[[#]] %[[Val]] %[[#]]
-; CHECK: %[[AggRes:.*]] = OpIAddCarry %[[Struct]] %[[PhiRes]] %[[Const1]]
-; CHECK: %[[Val]] = OpCompositeExtract %[[Int]] %[[AggRes]] 0
-; CHECK: %[[Over:.*]] = OpCompositeExtract %[[Int]] %[[AggRes]] 1
-; CHECK: %[[IsOver]] = OpINotEqual %[[Bool:.*]] %[[Over]] %[[Zero]]
-; CHECK: OpBranchConditional %[[IsOver]] %[[#]] %[[#]]
-; CHECK: OpStore %[[Ptr]] %[[Const42]] Aligned 1
+; CHECK: %[[APlusOne:.*]] = OpIAdd %[[Int]] %[[A]] %[[Const1]]
+; CHECK: OpBranch %[[#]]
+; CHECK: [[#]] = OpLabel
+; CHECK: %[[PhiRes]] = OpPhi %[[Int]] %[[Val]] %[[#]] %[[APlusOne]] %[[#]]
+; CHECK: %[[IsOver]] = OpIEqual %[[Bool]] %[[#]] %[[#]]
+; CHECK: OpBranchConditional %[[IsOver]] %[[#]] %[[#]]
+; CHECK: [[#]] = OpLabel
+; CHECK: OpStore %[[Ptr]] %[[Const42]] Aligned 1
+; CHECK: [[Val]] = OpIAdd %[[Int]] %[[PhiRes]] %[[Const1]]
; CHECK: OpBranch %[[#]]
-; CHECK: %[[#]] = OpLabel
-; CHECK: OpReturnValue %[[Val]]
-; CHECK: OpFunctionEnd
+; CHECK: [[#]] = OpLabel
+; OpReturnValue %[[PhiRes]]
+
+; NOLSR-DAG: OpName %[[Val:.*]] "math"
+; NOLSR-DAG: OpName %[[IsOver:.*]] "ov"
+; NOLSR-DAG: %[[Int:.*]] = OpTypeInt 32 0
+; NOLSR-DAG: %[[Char:.*]] = OpTypeInt 8 0
+; NOLSR-DAG: %[[PtrChar:.*]] = OpTypePointer Generic %[[Char]]
+; NOLSR-DAG: %[[Bool:.*]] = OpTypeBool
+; NOLSR-DAG: %[[Struct:.*]] = OpTypeStruct %[[Int]] %[[Int]]
+; NOLSR-DAG: %[[Const1:.*]] = OpConstant %[[Int]] 1
+; NOLSR-DAG: %[[Const42:.*]] = OpConstant %[[Char]] 42
+; NOLSR-DAG: %[[Zero:.*]] = OpConstantNull %[[Int]]
+
+; NOLSR: OpFunction
+; NOLSR: %[[A:.*]] = OpFunctionParameter %[[Int]]
+; NOLSR: %[[Ptr:.*]] = OpFunctionParameter %[[PtrChar]]
+; NOLSR: %[[#]] = OpLabel
+; NOLSR: OpBranch %[[#]]
+; NOLSR: %[[#]] = OpLabel
+; NOLSR: %[[PhiRes:.*]] = OpPhi %[[Int]] %[[A]] %[[#]] %[[Val]] %[[#]]
+; NOLSR: %[[AggRes:.*]] = OpIAddCarry %[[Struct]] %[[PhiRes]] %[[Const1]]
+; NOLSR: %[[Val]] = OpCompositeExtract %[[Int]] %[[AggRes]] 0
+; NOLSR: %[[Over:.*]] = OpCompositeExtract %[[Int]] %[[AggRes]] 1
+; NOLSR: %[[IsOver]] = OpINotEqual %[[Bool:.*]] %[[Over]] %[[Zero]]
+; NOLSR: OpBranchConditional %[[IsOver]] %[[#]] %[[#]]
+; NOLSR: OpStore %[[Ptr]] %[[Const42]] Aligned 1
+; NOLSR: OpBranch %[[#]]
+; NOLSR: %[[#]] = OpLabel
+; NOLSR: OpReturnValue %[[Val]]
+; NOLSR: OpFunctionEnd
define spir_func i32 @foo(i32 %a, ptr addrspace(4) %p) {
entry:
More information about the cfe-commits
mailing list