[llvm] [SPIR-V] Fix type mismatch in scalar-to-vector promotion for mixed-type builtins (PR #190969)
Arseniy Obolenskiy via llvm-commits
llvm-commits at lists.llvm.org
Wed Apr 8 06:20:31 PDT 2026
https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/190969
>From b24ea23a4f5bcd9bd1038ee588901e4ee454c6e9 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Wed, 8 Apr 2026 15:07:17 +0200
Subject: [PATCH 1/2] [SPIR-V] Fix type mismatch in scalar-to-vector promotion
for mixed-type builtins
When promoting scalar arguments to vectors for builtins like ldexp, pown, and rootn, use the correct vector type matching the argument element type instead of always using the return type
---
llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp | 10 ++++++++--
llvm/test/CodeGen/SPIRV/transcoding/ldexp.ll | 1 +
2 files changed, 9 insertions(+), 2 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
index db5218ec73bb7..88a43ec04f376 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
@@ -1186,6 +1186,9 @@ static bool builtinMayNeedPromotionToVec(uint32_t BuiltinNumber) {
case SPIRV::OpenCLExtInst::mix:
case SPIRV::OpenCLExtInst::step:
case SPIRV::OpenCLExtInst::smoothstep:
+ case SPIRV::OpenCLExtInst::ldexp:
+ case SPIRV::OpenCLExtInst::pown:
+ case SPIRV::OpenCLExtInst::rootn:
return true;
default:
break;
@@ -1215,10 +1218,13 @@ getBuiltinCallArguments(const SPIRV::IncomingCall *Call, uint32_t BuiltinNumber,
Register VecArg = Argument;
SPIRVTypeInst ArgumentType = GR->getSPIRVTypeForVReg(Argument);
if (ArgumentType != Call->ReturnType) {
- VecArg = createVirtualRegister(Call->ReturnType, GR, MIRBuilder);
+ SPIRVTypeInst VecType = GR->getOrCreateSPIRVVectorType(
+ ArgumentType, ResultElementCount, MIRBuilder, /*EmitIR=*/true);
+ VecArg = createVirtualRegister(VecType, GR, MIRBuilder);
+ Register VecTypeId = GR->getSPIRVTypeID(VecType);
auto VecSplat = MIRBuilder.buildInstr(SPIRV::OpCompositeConstruct)
.addDef(VecArg)
- .addUse(ReturnTypeId);
+ .addUse(VecTypeId);
for (unsigned I = 0; I != ResultElementCount; ++I)
VecSplat.addUse(Argument);
}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/ldexp.ll b/llvm/test/CodeGen/SPIRV/transcoding/ldexp.ll
index f1625e7665ad7..f5c4b274020ec 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/ldexp.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/ldexp.ll
@@ -1,6 +1,7 @@
;; Check that backend converts scalar arg to vector for ldexp math instructions
; 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 %}
;; #pragma OPENCL EXTENSION cl_khr_fp16 : enable
;; #pragma OPENCL EXTENSION cl_khr_fp64 : enable
>From 9ccd66cba80ac7ed5c789d5e46c57e2adca3724d Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Wed, 8 Apr 2026 15:13:18 +0200
Subject: [PATCH 2/2] Add pown and rootn tests
---
llvm/test/CodeGen/SPIRV/transcoding/pown.ll | 38 ++++++++++++++++++++
llvm/test/CodeGen/SPIRV/transcoding/rootn.ll | 38 ++++++++++++++++++++
2 files changed, 76 insertions(+)
create mode 100644 llvm/test/CodeGen/SPIRV/transcoding/pown.ll
create mode 100644 llvm/test/CodeGen/SPIRV/transcoding/rootn.ll
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/pown.ll b/llvm/test/CodeGen/SPIRV/transcoding/pown.ll
new file mode 100644
index 0000000000000..0f7509f0c1a70
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/pown.ll
@@ -0,0 +1,38 @@
+;; Check that backend converts scalar arg to vector for pown math instructions
+
+; 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 %}
+
+;; __kernel void test_kernel_float(float3 x, int n, __global float3* ret) {
+;; *ret = pown(x, n);
+;; }
+
+; CHECK-SPIRV: %{{.*}} pown
+
+define dso_local spir_kernel void @test_kernel_float(<3 x float> noundef %x, i32 noundef %n, ptr addrspace(1) nocapture noundef writeonly %ret) local_unnamed_addr {
+entry:
+ %call = call spir_func <3 x float> @_Z4pownDv3_fi(<3 x float> noundef %x, i32 noundef %n)
+ %extractVec2 = shufflevector <3 x float> %call, <3 x float> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 poison>
+ %storetmp3 = bitcast ptr addrspace(1) %ret to ptr addrspace(1)
+ store <4 x float> %extractVec2, ptr addrspace(1) %storetmp3, align 16
+ ret void
+}
+
+declare spir_func <3 x float> @_Z4pownDv3_fi(<3 x float> noundef, i32 noundef) local_unnamed_addr
+
+;; __kernel void test_kernel_double(double3 x, int n, __global double3* ret) {
+;; *ret = pown(x, n);
+;; }
+
+; CHECK-SPIRV: %{{.*}} pown
+
+define dso_local spir_kernel void @test_kernel_double(<3 x double> noundef %x, i32 noundef %n, ptr addrspace(1) nocapture noundef writeonly %ret) local_unnamed_addr {
+entry:
+ %call = call spir_func <3 x double> @_Z4pownDv3_di(<3 x double> noundef %x, i32 noundef %n)
+ %extractVec2 = shufflevector <3 x double> %call, <3 x double> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 poison>
+ %storetmp3 = bitcast ptr addrspace(1) %ret to ptr addrspace(1)
+ store <4 x double> %extractVec2, ptr addrspace(1) %storetmp3, align 32
+ ret void
+}
+
+declare spir_func <3 x double> @_Z4pownDv3_di(<3 x double> noundef, i32 noundef) local_unnamed_addr
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/rootn.ll b/llvm/test/CodeGen/SPIRV/transcoding/rootn.ll
new file mode 100644
index 0000000000000..92924fad8857a
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/rootn.ll
@@ -0,0 +1,38 @@
+;; Check that backend converts scalar arg to vector for rootn math instructions
+
+; 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 %}
+
+;; __kernel void test_kernel_float(float3 x, int n, __global float3* ret) {
+;; *ret = rootn(x, n);
+;; }
+
+; CHECK-SPIRV: %{{.*}} rootn
+
+define dso_local spir_kernel void @test_kernel_float(<3 x float> noundef %x, i32 noundef %n, ptr addrspace(1) nocapture noundef writeonly %ret) local_unnamed_addr {
+entry:
+ %call = call spir_func <3 x float> @_Z5rootnDv3_fi(<3 x float> noundef %x, i32 noundef %n)
+ %extractVec2 = shufflevector <3 x float> %call, <3 x float> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 poison>
+ %storetmp3 = bitcast ptr addrspace(1) %ret to ptr addrspace(1)
+ store <4 x float> %extractVec2, ptr addrspace(1) %storetmp3, align 16
+ ret void
+}
+
+declare spir_func <3 x float> @_Z5rootnDv3_fi(<3 x float> noundef, i32 noundef) local_unnamed_addr
+
+;; __kernel void test_kernel_double(double3 x, int n, __global double3* ret) {
+;; *ret = rootn(x, n);
+;; }
+
+; CHECK-SPIRV: %{{.*}} rootn
+
+define dso_local spir_kernel void @test_kernel_double(<3 x double> noundef %x, i32 noundef %n, ptr addrspace(1) nocapture noundef writeonly %ret) local_unnamed_addr {
+entry:
+ %call = call spir_func <3 x double> @_Z5rootnDv3_di(<3 x double> noundef %x, i32 noundef %n)
+ %extractVec2 = shufflevector <3 x double> %call, <3 x double> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 poison>
+ %storetmp3 = bitcast ptr addrspace(1) %ret to ptr addrspace(1)
+ store <4 x double> %extractVec2, ptr addrspace(1) %storetmp3, align 32
+ ret void
+}
+
+declare spir_func <3 x double> @_Z5rootnDv3_di(<3 x double> noundef, i32 noundef) local_unnamed_addr
More information about the llvm-commits
mailing list