llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-spir-v Author: Alex Voicu (AlexVlx) <details> <summary>Changes</summary> 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). As part of the update, we also delete a test that was fragile and depended on very specific CodeGenPrepare behaviour, around capability (integer ops with overflow intrisics) that already has coverage. --- Full diff: https://github.com/llvm/llvm-project/pull/110695.diff 5 Files Affected: - (modified) clang/lib/Basic/Targets/SPIR.h (+8-8) - (modified) clang/test/CodeGen/target-data.c (+1-1) - (modified) clang/test/CodeGenOpenCL/builtins-amdgcn.cl (+1-1) - (modified) llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp (+6-6) - (removed) llvm/test/CodeGen/SPIRV/optimizations/add-check-overflow.ll (-56) ``````````diff diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h index cc79562de2871e..09d4ad3c0ac620 100644 --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -314,8 +314,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, @@ -338,8 +338,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, @@ -362,8 +362,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, @@ -388,8 +388,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 8548aa00cfe877..fa875fe68b0c5b 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 e5384b2eb2c2c1..50c881a19cf58b 100644 --- a/llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp @@ -54,14 +54,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 deleted file mode 100644 index 1a630f77a44c5d..00000000000000 --- a/llvm/test/CodeGen/SPIRV/optimizations/add-check-overflow.ll +++ /dev/null @@ -1,56 +0,0 @@ -; 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). - -; 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 %} - -; 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" -; 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: %[[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: OpBranch %[[#]] -; CHECK: %[[#]] = OpLabel -; CHECK: OpReturnValue %[[Val]] -; CHECK: OpFunctionEnd - -define spir_func i32 @foo(i32 %a, ptr addrspace(4) %p) { -entry: - br label %l1 - -body: - store i8 42, ptr addrspace(4) %p - br label %l1 - -l1: - %e = phi i32 [ %a, %entry ], [ %i, %body ] - %i = add nsw i32 %e, 1 - %fl = icmp eq i32 %i, 0 - br i1 %fl, label %exit, label %body - -exit: - ret i32 %i -} `````````` </details> https://github.com/llvm/llvm-project/pull/110695 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
