[llvm] 3c2a9c9 - [SPIR-V] Fix type mismatch in scalar-to-vector promotion for mixed-type builtins (#190969)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Apr 13 06:26:11 PDT 2026
Author: Arseniy Obolenskiy
Date: 2026-04-13T15:26:05+02:00
New Revision: 3c2a9c90cb3c816609a6baabf919c392f0bcb40c
URL: https://github.com/llvm/llvm-project/commit/3c2a9c90cb3c816609a6baabf919c392f0bcb40c
DIFF: https://github.com/llvm/llvm-project/commit/3c2a9c90cb3c816609a6baabf919c392f0bcb40c.diff
LOG: [SPIR-V] Fix type mismatch in scalar-to-vector promotion for mixed-type builtins (#190969)
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: these builtins
take an integer argument but at the same time have floating point return
type
Fix `ldexp` test that does not pass spirv-val and add similar tests for
`pown` and `rootn`
related to https://github.com/llvm/llvm-project/issues/190736
Added:
llvm/test/CodeGen/SPIRV/transcoding/pown.ll
llvm/test/CodeGen/SPIRV/transcoding/rootn.ll
Modified:
llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
llvm/test/CodeGen/SPIRV/transcoding/ldexp.ll
Removed:
################################################################################
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
index db5218ec73bb7..b228799f77dac 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;
@@ -1214,11 +1217,15 @@ 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) {
- VecArg = createVirtualRegister(Call->ReturnType, GR, MIRBuilder);
+ if (GR->getScalarOrVectorComponentCount(ArgumentType) == 1 &&
+ ArgumentType != Call->ReturnType) {
+ 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..927eb5fa87495 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
+;; 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: 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
new file mode 100644
index 0000000000000..ff0ccd33dff70
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/pown.ll
@@ -0,0 +1,40 @@
+;; Check that backend converts scalar arg to vector for pown math instructions.
+
+; 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_fp64 : enable
+
+;; __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..1c72a5998e4b8
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/rootn.ll
@@ -0,0 +1,40 @@
+;; Check that backend converts scalar arg to vector for rootn math instructions.
+
+; 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_fp64 : enable
+
+;; __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