https://github.com/AlexVlx updated
https://github.com/llvm/llvm-project/pull/110695
>From 758fb6e28844d89031b5497d651cb2a9b71b6a0e Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.vo...@amd.com>
Date: Tue, 1 Oct 2024 17:10:50 +0100
Subject: [PATCH] Explicitly encode native integer widths for SPIR-V.
---
clang/lib/Basic/Targets/SPIR.h | 16 +++---
clang/test/CodeGen/target-data.c | 2 +-
clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 2 +-
llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp | 12 ++--
.../SPIRV/optimizations/add-check-overflow.ll | 56 -------------------
5 files changed, 16 insertions(+), 72 deletions(-)
delete mode 100644 llvm/test/CodeGen/SPIRV/optimizations/add-check-overflow.ll
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
-}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
