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: _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
