[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
Mon Apr 13 01:44:19 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/5] [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/5] 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
>From 5a9dcb79d114b3b30e01a14250bad5c0403705b0 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Wed, 8 Apr 2026 16:29:22 +0200
Subject: [PATCH 3/5] Address review comment
---
llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
index 88a43ec04f376..b228799f77dac 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
@@ -1217,7 +1217,8 @@ getBuiltinCallArguments(const SPIRV::IncomingCall *Call, uint32_t BuiltinNumber,
for (Register Argument : Call->Arguments) {
Register VecArg = Argument;
SPIRVTypeInst ArgumentType = GR->getSPIRVTypeForVReg(Argument);
- if (ArgumentType != Call->ReturnType) {
+ if (GR->getScalarOrVectorComponentCount(ArgumentType) == 1 &&
+ ArgumentType != Call->ReturnType) {
SPIRVTypeInst VecType = GR->getOrCreateSPIRVVectorType(
ArgumentType, ResultElementCount, MIRBuilder, /*EmitIR=*/true);
VecArg = createVirtualRegister(VecType, GR, MIRBuilder);
>From 28110520af6760c5905b8f3af942476117ae1227 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Mon, 13 Apr 2026 07:20:04 +0200
Subject: [PATCH 4/5] Address comments
---
llvm/test/CodeGen/SPIRV/transcoding/ldexp.ll | 2 +-
llvm/test/CodeGen/SPIRV/transcoding/pown.ll | 2 +-
llvm/test/CodeGen/SPIRV/transcoding/rootn.ll | 2 +-
3 files changed, 3 insertions(+), 3 deletions(-)
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/ldexp.ll b/llvm/test/CodeGen/SPIRV/transcoding/ldexp.ll
index f5c4b274020ec..5957e90c4e48c 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/ldexp.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/ldexp.ll
@@ -1,4 +1,4 @@
-;; Check that backend converts scalar arg to vector for ldexp math instructions
+;; 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 %}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/pown.ll b/llvm/test/CodeGen/SPIRV/transcoding/pown.ll
index 0f7509f0c1a70..38a14ccee24bc 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/pown.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/pown.ll
@@ -1,4 +1,4 @@
-;; Check that backend converts scalar arg to vector for pown math instructions
+;; 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 %}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/rootn.ll b/llvm/test/CodeGen/SPIRV/transcoding/rootn.ll
index 92924fad8857a..1d07455bac6c0 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/rootn.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/rootn.ll
@@ -1,4 +1,4 @@
-;; Check that backend converts scalar arg to vector for rootn math instructions
+;; 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 %}
>From 5ed9ac732482bc4cad7ebada441ee9566d994caf Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Mon, 13 Apr 2026 10:44:05 +0200
Subject: [PATCH 5/5] Add -verify-machineinstrs
---
llvm/test/CodeGen/SPIRV/transcoding/ldexp.ll | 4 ++--
llvm/test/CodeGen/SPIRV/transcoding/pown.ll | 4 ++--
llvm/test/CodeGen/SPIRV/transcoding/rootn.ll | 4 ++--
3 files changed, 6 insertions(+), 6 deletions(-)
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/ldexp.ll b/llvm/test/CodeGen/SPIRV/transcoding/ldexp.ll
index 5957e90c4e48c..927eb5fa87495 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/ldexp.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/ldexp.ll
@@ -1,7 +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 %}
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -verify-machineinstrs -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
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/pown.ll b/llvm/test/CodeGen/SPIRV/transcoding/pown.ll
index 38a14ccee24bc..5d63bb3b281e3 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/pown.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/pown.ll
@@ -1,7 +1,7 @@
;; 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 %}
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -verify-machineinstrs -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);
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/rootn.ll b/llvm/test/CodeGen/SPIRV/transcoding/rootn.ll
index 1d07455bac6c0..9a788b4c32122 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/rootn.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/rootn.ll
@@ -1,7 +1,7 @@
;; 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 %}
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -verify-machineinstrs -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);
More information about the llvm-commits
mailing list