[llvm] [SPIR-V] Add tests for lowering unmangled builtins calls (PR #84319)
Michal Paszkowski via llvm-commits
llvm-commits at lists.llvm.org
Thu Mar 7 19:17:37 PST 2024
https://github.com/michalpaszkowski updated https://github.com/llvm/llvm-project/pull/84319
>From b140785acec9492e3e88d660a2b3c8c51f7a61c4 Mon Sep 17 00:00:00 2001
From: Michal Paszkowski <michal at paszkowski.org>
Date: Thu, 7 Mar 2024 03:27:44 -0800
Subject: [PATCH 1/4] [SPIR-V] Move SPIR-V builtins tests to /builtins/
---
.../CodeGen/SPIRV/{ => builtins}/FOrdGreaterThanEqual_bool.ll | 0
llvm/test/CodeGen/SPIRV/{ => builtins}/SampledImageRetType.ll | 0
.../SPIRV/{transcoding => builtins}/SpecConstantComposite.ll | 0
.../SPIRV/{SpecConstants => builtins}/bool-spirv-specconstant.ll | 0
.../test/CodeGen/SPIRV/{transcoding => builtins}/builtin_calls.ll | 0
llvm/test/CodeGen/SPIRV/{ => builtins}/builtin_vars-decorate.ll | 0
llvm/test/CodeGen/SPIRV/{transcoding => builtins}/builtin_vars.ll | 0
.../SPIRV/{transcoding => builtins}/builtin_vars_arithmetics.ll | 0
.../CodeGen/SPIRV/{transcoding => builtins}/builtin_vars_opt.ll | 0
llvm/test/CodeGen/SPIRV/{ => builtins}/relationals.ll | 0
llvm/test/CodeGen/SPIRV/{ => builtins}/select.ll | 0
llvm/test/CodeGen/SPIRV/{transcoding => builtins}/spec_const.ll | 0
llvm/test/CodeGen/SPIRV/{ => builtins}/spec_const_decoration.ll | 0
llvm/test/CodeGen/SPIRV/{ => builtins}/spirv-load-store.ll | 0
llvm/test/CodeGen/SPIRV/{transcoding => builtins}/spirv-types.ll | 0
15 files changed, 0 insertions(+), 0 deletions(-)
rename llvm/test/CodeGen/SPIRV/{ => builtins}/FOrdGreaterThanEqual_bool.ll (100%)
rename llvm/test/CodeGen/SPIRV/{ => builtins}/SampledImageRetType.ll (100%)
rename llvm/test/CodeGen/SPIRV/{transcoding => builtins}/SpecConstantComposite.ll (100%)
rename llvm/test/CodeGen/SPIRV/{SpecConstants => builtins}/bool-spirv-specconstant.ll (100%)
rename llvm/test/CodeGen/SPIRV/{transcoding => builtins}/builtin_calls.ll (100%)
rename llvm/test/CodeGen/SPIRV/{ => builtins}/builtin_vars-decorate.ll (100%)
rename llvm/test/CodeGen/SPIRV/{transcoding => builtins}/builtin_vars.ll (100%)
rename llvm/test/CodeGen/SPIRV/{transcoding => builtins}/builtin_vars_arithmetics.ll (100%)
rename llvm/test/CodeGen/SPIRV/{transcoding => builtins}/builtin_vars_opt.ll (100%)
rename llvm/test/CodeGen/SPIRV/{ => builtins}/relationals.ll (100%)
rename llvm/test/CodeGen/SPIRV/{ => builtins}/select.ll (100%)
rename llvm/test/CodeGen/SPIRV/{transcoding => builtins}/spec_const.ll (100%)
rename llvm/test/CodeGen/SPIRV/{ => builtins}/spec_const_decoration.ll (100%)
rename llvm/test/CodeGen/SPIRV/{ => builtins}/spirv-load-store.ll (100%)
rename llvm/test/CodeGen/SPIRV/{transcoding => builtins}/spirv-types.ll (100%)
diff --git a/llvm/test/CodeGen/SPIRV/FOrdGreaterThanEqual_bool.ll b/llvm/test/CodeGen/SPIRV/builtins/FOrdGreaterThanEqual_bool.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/FOrdGreaterThanEqual_bool.ll
rename to llvm/test/CodeGen/SPIRV/builtins/FOrdGreaterThanEqual_bool.ll
diff --git a/llvm/test/CodeGen/SPIRV/SampledImageRetType.ll b/llvm/test/CodeGen/SPIRV/builtins/SampledImageRetType.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/SampledImageRetType.ll
rename to llvm/test/CodeGen/SPIRV/builtins/SampledImageRetType.ll
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/SpecConstantComposite.ll b/llvm/test/CodeGen/SPIRV/builtins/SpecConstantComposite.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/transcoding/SpecConstantComposite.ll
rename to llvm/test/CodeGen/SPIRV/builtins/SpecConstantComposite.ll
diff --git a/llvm/test/CodeGen/SPIRV/SpecConstants/bool-spirv-specconstant.ll b/llvm/test/CodeGen/SPIRV/builtins/bool-spirv-specconstant.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/SpecConstants/bool-spirv-specconstant.ll
rename to llvm/test/CodeGen/SPIRV/builtins/bool-spirv-specconstant.ll
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/builtin_calls.ll b/llvm/test/CodeGen/SPIRV/builtins/builtin_calls.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/transcoding/builtin_calls.ll
rename to llvm/test/CodeGen/SPIRV/builtins/builtin_calls.ll
diff --git a/llvm/test/CodeGen/SPIRV/builtin_vars-decorate.ll b/llvm/test/CodeGen/SPIRV/builtins/builtin_vars-decorate.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/builtin_vars-decorate.ll
rename to llvm/test/CodeGen/SPIRV/builtins/builtin_vars-decorate.ll
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars.ll b/llvm/test/CodeGen/SPIRV/builtins/builtin_vars.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/transcoding/builtin_vars.ll
rename to llvm/test/CodeGen/SPIRV/builtins/builtin_vars.ll
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_arithmetics.ll b/llvm/test/CodeGen/SPIRV/builtins/builtin_vars_arithmetics.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_arithmetics.ll
rename to llvm/test/CodeGen/SPIRV/builtins/builtin_vars_arithmetics.ll
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_opt.ll b/llvm/test/CodeGen/SPIRV/builtins/builtin_vars_opt.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_opt.ll
rename to llvm/test/CodeGen/SPIRV/builtins/builtin_vars_opt.ll
diff --git a/llvm/test/CodeGen/SPIRV/relationals.ll b/llvm/test/CodeGen/SPIRV/builtins/relationals.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/relationals.ll
rename to llvm/test/CodeGen/SPIRV/builtins/relationals.ll
diff --git a/llvm/test/CodeGen/SPIRV/select.ll b/llvm/test/CodeGen/SPIRV/builtins/select.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/select.ll
rename to llvm/test/CodeGen/SPIRV/builtins/select.ll
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/spec_const.ll b/llvm/test/CodeGen/SPIRV/builtins/spec_const.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/transcoding/spec_const.ll
rename to llvm/test/CodeGen/SPIRV/builtins/spec_const.ll
diff --git a/llvm/test/CodeGen/SPIRV/spec_const_decoration.ll b/llvm/test/CodeGen/SPIRV/builtins/spec_const_decoration.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/spec_const_decoration.ll
rename to llvm/test/CodeGen/SPIRV/builtins/spec_const_decoration.ll
diff --git a/llvm/test/CodeGen/SPIRV/spirv-load-store.ll b/llvm/test/CodeGen/SPIRV/builtins/spirv-load-store.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/spirv-load-store.ll
rename to llvm/test/CodeGen/SPIRV/builtins/spirv-load-store.ll
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/spirv-types.ll b/llvm/test/CodeGen/SPIRV/builtins/spirv-types.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/transcoding/spirv-types.ll
rename to llvm/test/CodeGen/SPIRV/builtins/spirv-types.ll
>From 000b488181064871dcd092c4bed9a615497214e2 Mon Sep 17 00:00:00 2001
From: Michal Paszkowski <michal at paszkowski.org>
Date: Thu, 7 Mar 2024 03:30:37 -0800
Subject: [PATCH 2/4] [SPIR-V] Rename SPIR-V builtin functions tests to match
the scheme
---
...{FOrdGreaterThanEqual_bool.ll => FOrdGreaterThanEqual-bool.ll} | 0
.../{SampledImageRetType.ll => SampledImage-return-type.ll} | 0
.../{spec_const_decoration.ll => SpecConstant-decoration.ll} | 0
.../CodeGen/SPIRV/builtins/{spec_const.ll => SpecConstant.ll} | 0
.../CodeGen/SPIRV/builtins/{builtin_calls.ll => builtin-calls.ll} | 0
.../{builtin_vars_arithmetics.ll => builtin-vars-arithmetics.ll} | 0
.../{builtin_vars-decorate.ll => builtin-vars-decorate.ll} | 0
.../SPIRV/builtins/{builtin_vars_opt.ll => builtin-vars-opt.ll} | 0
.../CodeGen/SPIRV/builtins/{builtin_vars.ll => builtin-vars.ll} | 0
9 files changed, 0 insertions(+), 0 deletions(-)
rename llvm/test/CodeGen/SPIRV/builtins/{FOrdGreaterThanEqual_bool.ll => FOrdGreaterThanEqual-bool.ll} (100%)
rename llvm/test/CodeGen/SPIRV/builtins/{SampledImageRetType.ll => SampledImage-return-type.ll} (100%)
rename llvm/test/CodeGen/SPIRV/builtins/{spec_const_decoration.ll => SpecConstant-decoration.ll} (100%)
rename llvm/test/CodeGen/SPIRV/builtins/{spec_const.ll => SpecConstant.ll} (100%)
rename llvm/test/CodeGen/SPIRV/builtins/{builtin_calls.ll => builtin-calls.ll} (100%)
rename llvm/test/CodeGen/SPIRV/builtins/{builtin_vars_arithmetics.ll => builtin-vars-arithmetics.ll} (100%)
rename llvm/test/CodeGen/SPIRV/builtins/{builtin_vars-decorate.ll => builtin-vars-decorate.ll} (100%)
rename llvm/test/CodeGen/SPIRV/builtins/{builtin_vars_opt.ll => builtin-vars-opt.ll} (100%)
rename llvm/test/CodeGen/SPIRV/builtins/{builtin_vars.ll => builtin-vars.ll} (100%)
diff --git a/llvm/test/CodeGen/SPIRV/builtins/FOrdGreaterThanEqual_bool.ll b/llvm/test/CodeGen/SPIRV/builtins/FOrdGreaterThanEqual-bool.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/builtins/FOrdGreaterThanEqual_bool.ll
rename to llvm/test/CodeGen/SPIRV/builtins/FOrdGreaterThanEqual-bool.ll
diff --git a/llvm/test/CodeGen/SPIRV/builtins/SampledImageRetType.ll b/llvm/test/CodeGen/SPIRV/builtins/SampledImage-return-type.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/builtins/SampledImageRetType.ll
rename to llvm/test/CodeGen/SPIRV/builtins/SampledImage-return-type.ll
diff --git a/llvm/test/CodeGen/SPIRV/builtins/spec_const_decoration.ll b/llvm/test/CodeGen/SPIRV/builtins/SpecConstant-decoration.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/builtins/spec_const_decoration.ll
rename to llvm/test/CodeGen/SPIRV/builtins/SpecConstant-decoration.ll
diff --git a/llvm/test/CodeGen/SPIRV/builtins/spec_const.ll b/llvm/test/CodeGen/SPIRV/builtins/SpecConstant.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/builtins/spec_const.ll
rename to llvm/test/CodeGen/SPIRV/builtins/SpecConstant.ll
diff --git a/llvm/test/CodeGen/SPIRV/builtins/builtin_calls.ll b/llvm/test/CodeGen/SPIRV/builtins/builtin-calls.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/builtins/builtin_calls.ll
rename to llvm/test/CodeGen/SPIRV/builtins/builtin-calls.ll
diff --git a/llvm/test/CodeGen/SPIRV/builtins/builtin_vars_arithmetics.ll b/llvm/test/CodeGen/SPIRV/builtins/builtin-vars-arithmetics.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/builtins/builtin_vars_arithmetics.ll
rename to llvm/test/CodeGen/SPIRV/builtins/builtin-vars-arithmetics.ll
diff --git a/llvm/test/CodeGen/SPIRV/builtins/builtin_vars-decorate.ll b/llvm/test/CodeGen/SPIRV/builtins/builtin-vars-decorate.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/builtins/builtin_vars-decorate.ll
rename to llvm/test/CodeGen/SPIRV/builtins/builtin-vars-decorate.ll
diff --git a/llvm/test/CodeGen/SPIRV/builtins/builtin_vars_opt.ll b/llvm/test/CodeGen/SPIRV/builtins/builtin-vars-opt.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/builtins/builtin_vars_opt.ll
rename to llvm/test/CodeGen/SPIRV/builtins/builtin-vars-opt.ll
diff --git a/llvm/test/CodeGen/SPIRV/builtins/builtin_vars.ll b/llvm/test/CodeGen/SPIRV/builtins/builtin-vars.ll
similarity index 100%
rename from llvm/test/CodeGen/SPIRV/builtins/builtin_vars.ll
rename to llvm/test/CodeGen/SPIRV/builtins/builtin-vars.ll
>From fa5b3524e6a003642a3e4a079e387ec8b75bcdc7 Mon Sep 17 00:00:00 2001
From: Michal Paszkowski <michal at paszkowski.org>
Date: Thu, 7 Mar 2024 05:07:05 -0800
Subject: [PATCH 3/4] [SPIR-V] Add tests for lowering unmangled builtins calls
---
.../builtins/FOrdGreaterThanEqual-bool.ll | 1 +
.../CodeGen/SPIRV/builtins/builtin-calls.ll | 1 +
.../CodeGen/SPIRV/builtins/relationals.ll | 1 +
llvm/test/CodeGen/SPIRV/builtins/select.ll | 1 +
.../SPIRV/builtins/spirv-load-store.ll | 10 +++++--
.../unmangled-SampledImage-return-type.ll | 29 +++++++++++++++++++
.../SPIRV/builtins/unmangled-builtin-calls.ll | 17 +++++++++++
.../SPIRV/builtins/unmangled-select.ll | 15 ++++++++++
.../builtins/unmangled-spirv-load-store.ll | 20 +++++++++++++
9 files changed, 92 insertions(+), 3 deletions(-)
create mode 100644 llvm/test/CodeGen/SPIRV/builtins/unmangled-SampledImage-return-type.ll
create mode 100644 llvm/test/CodeGen/SPIRV/builtins/unmangled-builtin-calls.ll
create mode 100644 llvm/test/CodeGen/SPIRV/builtins/unmangled-select.ll
create mode 100644 llvm/test/CodeGen/SPIRV/builtins/unmangled-spirv-load-store.ll
diff --git a/llvm/test/CodeGen/SPIRV/builtins/FOrdGreaterThanEqual-bool.ll b/llvm/test/CodeGen/SPIRV/builtins/FOrdGreaterThanEqual-bool.ll
index 735b35d757e7d8..24491f313490b2 100644
--- a/llvm/test/CodeGen/SPIRV/builtins/FOrdGreaterThanEqual-bool.ll
+++ b/llvm/test/CodeGen/SPIRV/builtins/FOrdGreaterThanEqual-bool.ll
@@ -1,4 +1,5 @@
; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
; CHECK-SPIRV: OpFOrdGreaterThanEqual
; CHECK-SPIRV-NOT: OpSelect
diff --git a/llvm/test/CodeGen/SPIRV/builtins/builtin-calls.ll b/llvm/test/CodeGen/SPIRV/builtins/builtin-calls.ll
index 9b1ce76631809d..0a02a8bf56ace5 100644
--- a/llvm/test/CodeGen/SPIRV/builtins/builtin-calls.ll
+++ b/llvm/test/CodeGen/SPIRV/builtins/builtin-calls.ll
@@ -1,4 +1,5 @@
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
; CHECK-SPIRV-DAG: OpDecorate %[[#Id:]] BuiltIn GlobalInvocationId
; CHECK-SPIRV-DAG: OpDecorate %[[#Id:]] BuiltIn GlobalLinearId
diff --git a/llvm/test/CodeGen/SPIRV/builtins/relationals.ll b/llvm/test/CodeGen/SPIRV/builtins/relationals.ll
index 1644dc7c03d911..13a9869ec17ed3 100644
--- a/llvm/test/CodeGen/SPIRV/builtins/relationals.ll
+++ b/llvm/test/CodeGen/SPIRV/builtins/relationals.ll
@@ -1,4 +1,5 @@
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
declare dso_local spir_func <4 x i8> @_Z13__spirv_IsNanIDv4_aDv4_fET_T0_(<4 x float>)
declare dso_local spir_func <4 x i8> @_Z13__spirv_IsInfIDv4_aDv4_fET_T0_(<4 x float>)
diff --git a/llvm/test/CodeGen/SPIRV/builtins/select.ll b/llvm/test/CodeGen/SPIRV/builtins/select.ll
index b34e91be1dbcda..6a42b63dcd34a6 100644
--- a/llvm/test/CodeGen/SPIRV/builtins/select.ll
+++ b/llvm/test/CodeGen/SPIRV/builtins/select.ll
@@ -1,4 +1,5 @@
; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
; CHECK-SPIRV: OpSelect
diff --git a/llvm/test/CodeGen/SPIRV/builtins/spirv-load-store.ll b/llvm/test/CodeGen/SPIRV/builtins/spirv-load-store.ll
index a82bf0ab2e01f6..8b80e0308dd0e1 100644
--- a/llvm/test/CodeGen/SPIRV/builtins/spirv-load-store.ll
+++ b/llvm/test/CodeGen/SPIRV/builtins/spirv-load-store.ll
@@ -1,9 +1,13 @@
; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
;; Translate SPIR-V friendly OpLoad and OpStore calls
-; CHECK: %[[#CONST:]] = OpConstant %[[#]] 42
-; CHECK: OpStore %[[#PTR:]] %[[#CONST]] Volatile|Aligned 4
-; CHECK: %[[#]] = OpLoad %[[#]] %[[#PTR]]
+; CHECK-DAG: %[[#INT32:]] = OpTypeInt 32 0
+; CHECK-DAG: %[[#CONST:]] = OpConstant %[[#INT32]] 42
+; CHECK-DAG: %[[#PTRINT32:]] = OpTypePointer CrossWorkgroup %[[#INT32]]
+; CHECK: %[[#BITCASTorPARAMETER:]] = {{OpBitcast|OpFunctionParameter}}{{.*}}%[[#PTRINT32]]{{.*}}
+
+; CHECK: OpStore %[[#BITCASTorPARAMETER]] %[[#CONST]] Volatile|Aligned 4
+; CHECK: %[[#]] = OpLoad %[[#]] %[[#BITCASTorPARAMETER]]
define weak_odr dso_local spir_kernel void @foo(i32 addrspace(1)* %var) {
entry:
diff --git a/llvm/test/CodeGen/SPIRV/builtins/unmangled-SampledImage-return-type.ll b/llvm/test/CodeGen/SPIRV/builtins/unmangled-SampledImage-return-type.ll
new file mode 100644
index 00000000000000..3459598735e65a
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/builtins/unmangled-SampledImage-return-type.ll
@@ -0,0 +1,29 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+; XFAIL: *
+
+; CHECK: %[[#image1d_t:]] = OpTypeImage
+; CHECK: %[[#sampler_t:]] = OpTypeSampler
+; CHECK: %[[#sampled_image_t:]] = OpTypeSampledImage
+
+declare dso_local spir_func ptr addrspace(4) @__spirv_SampledImage(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) %0, target("spirv.Sampler") %1) local_unnamed_addr
+
+declare dso_local spir_func <4 x float> @__spirv_ImageSampleExplicitLod(ptr addrspace(4) %0, i32 %1, i32 %2, float %3) local_unnamed_addr
+
+ at __spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(2) constant <3 x i64>, align 32
+
+define weak_odr dso_local spir_kernel void @_ZTS17image_kernel_readILi1EE(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0), target("spirv.Sampler")) {
+; CHECK: OpFunction
+; CHECK: %[[#image:]] = OpFunctionParameter %[[#image1d_t]]
+; CHECK: %[[#sampler:]] = OpFunctionParameter %[[#sampler_t]]
+ %3 = load <3 x i64>, ptr addrspace(2) @__spirv_BuiltInGlobalInvocationId, align 32
+ %4 = extractelement <3 x i64> %3, i64 0
+ %5 = trunc i64 %4 to i32
+ %6 = call spir_func ptr addrspace(4) @__spirv_SampledImage(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) %0, target("spirv.Sampler") %1)
+ %7 = call spir_func <4 x float> @__spirv_ImageSampleExplicitLod(ptr addrspace(4) %6, i32 %5, i32 2, float 0.000000e+00)
+
+; CHECK: %[[#sampled_image:]] = OpSampledImage %[[#sampled_image_t]] %[[#image]] %[[#sampler]]
+; CHECK: %[[#]] = OpImageSampleExplicitLod %[[#]] %[[#sampled_image]] %[[#]] {{.*}} %[[#]]
+
+ ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/builtins/unmangled-builtin-calls.ll b/llvm/test/CodeGen/SPIRV/builtins/unmangled-builtin-calls.ll
new file mode 100644
index 00000000000000..4519d52dc8f30e
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/builtins/unmangled-builtin-calls.ll
@@ -0,0 +1,17 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; CHECK-SPIRV-DAG: OpDecorate %[[#Id:]] BuiltIn GlobalInvocationId
+; CHECK-SPIRV-DAG: OpDecorate %[[#Id:]] BuiltIn GlobalLinearId
+; CHECK-SPIRV: %[[#Id:]] = OpVariable %[[#]]
+; CHECK-SPIRV: %[[#Id:]] = OpVariable %[[#]]
+
+define spir_kernel void @f() {
+entry:
+ %0 = call spir_func i32 @__spirv_BuiltInGlobalLinearId()
+ %1 = call spir_func i64 @__spirv_BuiltInGlobalInvocationId(i32 1)
+ ret void
+}
+
+declare spir_func i32 @__spirv_BuiltInGlobalLinearId()
+declare spir_func i64 @__spirv_BuiltInGlobalInvocationId(i32)
diff --git a/llvm/test/CodeGen/SPIRV/builtins/unmangled-select.ll b/llvm/test/CodeGen/SPIRV/builtins/unmangled-select.ll
new file mode 100644
index 00000000000000..e88fa49a8f55f9
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/builtins/unmangled-select.ll
@@ -0,0 +1,15 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; CHECK-SPIRV: OpSelect
+
+;; LLVM IR was generated with -cl-std=c++ option
+
+define spir_kernel void @test(i32 %op1, i32 %op2) {
+entry:
+ %0 = trunc i8 undef to i1
+ %call = call spir_func i32 @__spirv_Select(i1 zeroext %0, i32 %op1, i32 %op2)
+ ret void
+}
+
+declare spir_func i32 @__spirv_Select(i1 zeroext, i32, i32)
diff --git a/llvm/test/CodeGen/SPIRV/builtins/unmangled-spirv-load-store.ll b/llvm/test/CodeGen/SPIRV/builtins/unmangled-spirv-load-store.ll
new file mode 100644
index 00000000000000..19fc7967a3b63b
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/builtins/unmangled-spirv-load-store.ll
@@ -0,0 +1,20 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK-DAG: %[[#INT8:]] = OpTypeInt 8 0
+; CHECK-DAG: %[[#INT32:]] = OpTypeInt 32 0
+; CHECK-DAG: %[[#CONST:]] = OpConstant %[[#INT32]] 42
+; CHECK-DAG: %[[#PTRINT8:]] = OpTypePointer CrossWorkgroup %[[#INT8]]
+; CHECK: %[[#BITCASTorPARAMETER:]] = {{OpBitcast|OpFunctionParameter}}{{.*}}%[[#PTRINT8]]{{.*}}
+
+; CHECK: OpStore %[[#BITCASTorPARAMETER]] %[[#CONST]] Volatile|Aligned 4
+; CHECK: %[[#]] = OpLoad %[[#]] %[[#BITCASTorPARAMETER]]
+
+define weak_odr dso_local spir_kernel void @foo(ptr addrspace(1) %var) {
+entry:
+ tail call spir_func void @__spirv_Store(ptr addrspace(1) %var, i32 42, i32 3, i32 4)
+ %value = tail call spir_func double @__spirv_Load(ptr addrspace(1) %var)
+ ret void
+}
+
+declare dso_local spir_func double @__spirv_Load(ptr addrspace(1)) local_unnamed_addr
+declare dso_local spir_func void @__spirv_Store(ptr addrspace(1), i32, i32, i32) local_unnamed_addr
>From 4b68f097503d583e3ac88e995e2dd9b19aa8872e Mon Sep 17 00:00:00 2001
From: Michal Paszkowski <michal at paszkowski.org>
Date: Thu, 7 Mar 2024 19:16:47 -0800
Subject: [PATCH 4/4] [SPIR-V] Renamed virtual registers in new tests
---
.../SPIRV/builtins/bool-spirv-specconstant.ll | 20 +++++++++----------
.../CodeGen/SPIRV/builtins/builtin-calls.ll | 4 ++--
.../CodeGen/SPIRV/builtins/relationals.ll | 12 +++++------
llvm/test/CodeGen/SPIRV/builtins/select.ll | 4 ++--
.../CodeGen/SPIRV/builtins/spirv-types.ll | 4 ++--
.../unmangled-SampledImage-return-type.ll | 12 +++++------
.../SPIRV/builtins/unmangled-builtin-calls.ll | 4 ++--
.../SPIRV/builtins/unmangled-select.ll | 4 ++--
8 files changed, 32 insertions(+), 32 deletions(-)
diff --git a/llvm/test/CodeGen/SPIRV/builtins/bool-spirv-specconstant.ll b/llvm/test/CodeGen/SPIRV/builtins/bool-spirv-specconstant.ll
index 6e414f79bdde54..75a3994f2bea01 100644
--- a/llvm/test/CodeGen/SPIRV/builtins/bool-spirv-specconstant.ll
+++ b/llvm/test/CodeGen/SPIRV/builtins/bool-spirv-specconstant.ll
@@ -15,16 +15,16 @@ $"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel1" = comdat any
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel1"(i8 addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3) local_unnamed_addr comdat {
entry:
- %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
- %1 = addrspacecast i64* %0 to i64 addrspace(4)*
- %2 = load i64, i64 addrspace(4)* %1, align 8
- %add.ptr.i = getelementptr inbounds i8, i8 addrspace(1)* %_arg_, i64 %2
- %3 = call i1 @_Z20__spirv_SpecConstantia(i32 0, i8 1)
- %ptridx.ascast.i.i = addrspacecast i8 addrspace(1)* %add.ptr.i to i8 addrspace(4)*
- %selected = select i1 %3, i8 0, i8 1
- %frombool.i = zext i1 %3 to i8
- %sum = add i8 %frombool.i, %selected
- store i8 %selected, i8 addrspace(4)* %ptridx.ascast.i.i, align 1
+ %gep1 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
+ %cast1 = addrspacecast i64* %gep1 to i64 addrspace(4)*
+ %load = load i64, i64 addrspace(4)* %cast1, align 8
+ %gep2 = getelementptr inbounds i8, i8 addrspace(1)* %_arg_, i64 %load
+ %call1 = call i1 @_Z20__spirv_SpecConstantia(i32 0, i8 1)
+ %cast2 = addrspacecast i8 addrspace(1)* %gep2 to i8 addrspace(4)*
+ %selected = select i1 %call1, i8 0, i8 1
+ %bool = zext i1 %call1 to i8
+ %sum = add i8 %bool, %selected
+ store i8 %selected, i8 addrspace(4)* %cast2, align 1
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/builtins/builtin-calls.ll b/llvm/test/CodeGen/SPIRV/builtins/builtin-calls.ll
index 0a02a8bf56ace5..0ba25d688952c4 100644
--- a/llvm/test/CodeGen/SPIRV/builtins/builtin-calls.ll
+++ b/llvm/test/CodeGen/SPIRV/builtins/builtin-calls.ll
@@ -8,8 +8,8 @@
define spir_kernel void @f() {
entry:
- %0 = call spir_func i32 @_Z29__spirv_BuiltInGlobalLinearIdv()
- %1 = call spir_func i64 @_Z33__spirv_BuiltInGlobalInvocationIdi(i32 1)
+ %call1 = call spir_func i32 @_Z29__spirv_BuiltInGlobalLinearIdv()
+ %call2 = call spir_func i64 @_Z33__spirv_BuiltInGlobalInvocationIdi(i32 1)
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/builtins/relationals.ll b/llvm/test/CodeGen/SPIRV/builtins/relationals.ll
index 13a9869ec17ed3..e6693ac96a5e60 100644
--- a/llvm/test/CodeGen/SPIRV/builtins/relationals.ll
+++ b/llvm/test/CodeGen/SPIRV/builtins/relationals.ll
@@ -14,28 +14,28 @@ define spir_kernel void @k() {
entry:
%arg1 = alloca <4 x float>, align 16
%ret = alloca <4 x i8>, align 4
- %0 = load <4 x float>, <4 x float>* %arg1, align 16
- %call1 = call spir_func <4 x i8> @_Z13__spirv_IsNanIDv4_aDv4_fET_T0_(<4 x float> %0)
+ %loaded = load <4 x float>, <4 x float>* %arg1, align 16
+ %call1 = call spir_func <4 x i8> @_Z13__spirv_IsNanIDv4_aDv4_fET_T0_(<4 x float> %loaded)
; CHECK-SPIRV: %[[#IsNanRes:]] = OpIsNan %[[#TBoolVec]]
; CHECK-SPIRV: %[[#SelectRes:]] = OpSelect %[[#]] %[[#IsNanRes]]
; CHECK-SPIRV: OpStore %[[#]] %[[#SelectRes]]
store <4 x i8> %call1, <4 x i8>* %ret, align 4
- %call2 = call spir_func <4 x i8> @_Z13__spirv_IsInfIDv4_aDv4_fET_T0_(<4 x float> %0)
+ %call2 = call spir_func <4 x i8> @_Z13__spirv_IsInfIDv4_aDv4_fET_T0_(<4 x float> %loaded)
; CHECK-SPIRV: %[[#IsInfRes:]] = OpIsInf %[[#TBoolVec]]
; CHECK-SPIRV: %[[#Select1Res:]] = OpSelect %[[#]] %[[#IsInfRes]]
; CHECK-SPIRV: OpStore %[[#]] %[[#Select1Res]]
store <4 x i8> %call2, <4 x i8>* %ret, align 4
- %call3 = call spir_func <4 x i8> @_Z16__spirv_IsFiniteIDv4_aDv4_fET_T0_(<4 x float> %0)
+ %call3 = call spir_func <4 x i8> @_Z16__spirv_IsFiniteIDv4_aDv4_fET_T0_(<4 x float> %loaded)
; CHECK-SPIRV: %[[#IsFiniteRes:]] = OpIsFinite %[[#TBoolVec]]
; CHECK-SPIRV: %[[#Select2Res:]] = OpSelect %[[#]] %[[#IsFiniteRes]]
; CHECK-SPIRV: OpStore %[[#]] %[[#Select2Res]]
store <4 x i8> %call3, <4 x i8>* %ret, align 4
- %call4 = call spir_func <4 x i8> @_Z16__spirv_IsNormalIDv4_aDv4_fET_T0_(<4 x float> %0)
+ %call4 = call spir_func <4 x i8> @_Z16__spirv_IsNormalIDv4_aDv4_fET_T0_(<4 x float> %loaded)
; CHECK-SPIRV: %[[#IsNormalRes:]] = OpIsNormal %[[#TBoolVec]]
; CHECK-SPIRV: %[[#Select3Res:]] = OpSelect %[[#]] %[[#IsNormalRes]]
; CHECK-SPIRV: OpStore %[[#]] %[[#Select3Res]]
store <4 x i8> %call4, <4 x i8>* %ret, align 4
- %call5 = call spir_func <4 x i8> @_Z18__spirv_SignBitSetIDv4_aDv4_fET_T0_(<4 x float> %0)
+ %call5 = call spir_func <4 x i8> @_Z18__spirv_SignBitSetIDv4_aDv4_fET_T0_(<4 x float> %loaded)
; CHECK-SPIRV: %[[#SignBitSetRes:]] = OpSignBitSet %[[#TBoolVec]]
; CHECK-SPIRV: %[[#Select4Res:]] = OpSelect %[[#]] %[[#SignBitSetRes]]
; CHECK-SPIRV: OpStore %[[#]] %[[#Select4Res]]
diff --git a/llvm/test/CodeGen/SPIRV/builtins/select.ll b/llvm/test/CodeGen/SPIRV/builtins/select.ll
index 6a42b63dcd34a6..d9ce6afaf0243e 100644
--- a/llvm/test/CodeGen/SPIRV/builtins/select.ll
+++ b/llvm/test/CodeGen/SPIRV/builtins/select.ll
@@ -7,8 +7,8 @@
define spir_kernel void @test(i32 %op1, i32 %op2) {
entry:
- %0 = trunc i8 undef to i1
- %call = call spir_func i32 @_Z14__spirv_Selectbii(i1 zeroext %0, i32 %op1, i32 %op2)
+ %truncated = trunc i8 undef to i1
+ %call = call spir_func i32 @_Z14__spirv_Selectbii(i1 zeroext %truncated, i32 %op1, i32 %op2)
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/builtins/spirv-types.ll b/llvm/test/CodeGen/SPIRV/builtins/spirv-types.ll
index f8db0ba2c1cd4e..1f8f56af7bef8c 100644
--- a/llvm/test/CodeGen/SPIRV/builtins/spirv-types.ll
+++ b/llvm/test/CodeGen/SPIRV/builtins/spirv-types.ll
@@ -75,8 +75,8 @@ define spir_func void @bar(
define spir_func void @test_sampler(target("spirv.Image", float, 1, 1, 0, 0, 0, 0, 0) %srcimg.coerce,
target("spirv.Sampler") %s.coerce) {
- %1 = tail call spir_func target("spirv.Image", float, 1, 1, 0, 0, 0, 0, 0) @_Z20__spirv_SampledImagePU3AS1K34__spirv_Image__float_1_1_0_0_0_0_0PU3AS1K15__spirv_Sampler(target("spirv.Image", float, 1, 1, 0, 0, 0, 0, 0) %srcimg.coerce, target("spirv.Sampler") %s.coerce)
- %2 = tail call spir_func <4 x float> @_Z38__spirv_ImageSampleExplicitLod_Rfloat4PU3AS120__spirv_SampledImageDv4_iif(target("spirv.Image", float, 1, 1, 0, 0, 0, 0, 0) %1, <4 x i32> zeroinitializer, i32 2, float 1.000000e+00)
+ %call1 = tail call spir_func target("spirv.Image", float, 1, 1, 0, 0, 0, 0, 0) @_Z20__spirv_SampledImagePU3AS1K34__spirv_Image__float_1_1_0_0_0_0_0PU3AS1K15__spirv_Sampler(target("spirv.Image", float, 1, 1, 0, 0, 0, 0, 0) %srcimg.coerce, target("spirv.Sampler") %s.coerce)
+ %call2 = tail call spir_func <4 x float> @_Z38__spirv_ImageSampleExplicitLod_Rfloat4PU3AS120__spirv_SampledImageDv4_iif(target("spirv.Image", float, 1, 1, 0, 0, 0, 0, 0) %call1, <4 x i32> zeroinitializer, i32 2, float 1.000000e+00)
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/builtins/unmangled-SampledImage-return-type.ll b/llvm/test/CodeGen/SPIRV/builtins/unmangled-SampledImage-return-type.ll
index 3459598735e65a..522d755f70c2ea 100644
--- a/llvm/test/CodeGen/SPIRV/builtins/unmangled-SampledImage-return-type.ll
+++ b/llvm/test/CodeGen/SPIRV/builtins/unmangled-SampledImage-return-type.ll
@@ -8,7 +8,7 @@
declare dso_local spir_func ptr addrspace(4) @__spirv_SampledImage(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) %0, target("spirv.Sampler") %1) local_unnamed_addr
-declare dso_local spir_func <4 x float> @__spirv_ImageSampleExplicitLod(ptr addrspace(4) %0, i32 %1, i32 %2, float %3) local_unnamed_addr
+declare dso_local spir_func <4 x float> @__spirv_ImageSampleExplicitLod(ptr addrspace(4) %0, i32 %1, i32 %2, float %loaded) local_unnamed_addr
@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(2) constant <3 x i64>, align 32
@@ -16,11 +16,11 @@ define weak_odr dso_local spir_kernel void @_ZTS17image_kernel_readILi1EE(target
; CHECK: OpFunction
; CHECK: %[[#image:]] = OpFunctionParameter %[[#image1d_t]]
; CHECK: %[[#sampler:]] = OpFunctionParameter %[[#sampler_t]]
- %3 = load <3 x i64>, ptr addrspace(2) @__spirv_BuiltInGlobalInvocationId, align 32
- %4 = extractelement <3 x i64> %3, i64 0
- %5 = trunc i64 %4 to i32
- %6 = call spir_func ptr addrspace(4) @__spirv_SampledImage(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) %0, target("spirv.Sampler") %1)
- %7 = call spir_func <4 x float> @__spirv_ImageSampleExplicitLod(ptr addrspace(4) %6, i32 %5, i32 2, float 0.000000e+00)
+ %loaded = load <3 x i64>, ptr addrspace(2) @__spirv_BuiltInGlobalInvocationId, align 32
+ %extracted = extractelement <3 x i64> %loaded, i64 0
+ %truncated = trunc i64 %extracted to i32
+ %call1 = call spir_func ptr addrspace(4) @__spirv_SampledImage(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) %0, target("spirv.Sampler") %1)
+ %call2 = call spir_func <4 x float> @__spirv_ImageSampleExplicitLod(ptr addrspace(4) %call1, i32 %truncated, i32 2, float 0.000000e+00)
; CHECK: %[[#sampled_image:]] = OpSampledImage %[[#sampled_image_t]] %[[#image]] %[[#sampler]]
; CHECK: %[[#]] = OpImageSampleExplicitLod %[[#]] %[[#sampled_image]] %[[#]] {{.*}} %[[#]]
diff --git a/llvm/test/CodeGen/SPIRV/builtins/unmangled-builtin-calls.ll b/llvm/test/CodeGen/SPIRV/builtins/unmangled-builtin-calls.ll
index 4519d52dc8f30e..1c1ffa03032993 100644
--- a/llvm/test/CodeGen/SPIRV/builtins/unmangled-builtin-calls.ll
+++ b/llvm/test/CodeGen/SPIRV/builtins/unmangled-builtin-calls.ll
@@ -8,8 +8,8 @@
define spir_kernel void @f() {
entry:
- %0 = call spir_func i32 @__spirv_BuiltInGlobalLinearId()
- %1 = call spir_func i64 @__spirv_BuiltInGlobalInvocationId(i32 1)
+ %call1 = call spir_func i32 @__spirv_BuiltInGlobalLinearId()
+ %call2 = call spir_func i64 @__spirv_BuiltInGlobalInvocationId(i32 1)
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/builtins/unmangled-select.ll b/llvm/test/CodeGen/SPIRV/builtins/unmangled-select.ll
index e88fa49a8f55f9..d6f04d41336c9f 100644
--- a/llvm/test/CodeGen/SPIRV/builtins/unmangled-select.ll
+++ b/llvm/test/CodeGen/SPIRV/builtins/unmangled-select.ll
@@ -7,8 +7,8 @@
define spir_kernel void @test(i32 %op1, i32 %op2) {
entry:
- %0 = trunc i8 undef to i1
- %call = call spir_func i32 @__spirv_Select(i1 zeroext %0, i32 %op1, i32 %op2)
+ %truncated = trunc i8 undef to i1
+ %call = call spir_func i32 @__spirv_Select(i1 zeroext %truncated, i32 %op1, i32 %op2)
ret void
}
More information about the llvm-commits
mailing list