[spirv-reader][ir] Implement OpConvertUToF
Implement the `OpConvertUToF` operand. This needs to make sure that the
argument is provided as an unsigned value. Bitcast as needed.
Bug: 391486027
Change-Id: Ib76643d09c649c8b42283af1cff8140d9a258484
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/226679
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/spirv/builtin_fn.cc b/src/tint/lang/spirv/builtin_fn.cc
index f70d75d..60cb6b9 100644
--- a/src/tint/lang/spirv/builtin_fn.cc
+++ b/src/tint/lang/spirv/builtin_fn.cc
@@ -168,6 +168,8 @@
return "convertFToS";
case BuiltinFn::kConvertSToF:
return "convertSToF";
+ case BuiltinFn::kConvertUToF:
+ return "convertUToF";
case BuiltinFn::kSdot:
return "sdot";
case BuiltinFn::kUdot:
@@ -260,6 +262,7 @@
case BuiltinFn::kSMod:
case BuiltinFn::kConvertFToS:
case BuiltinFn::kConvertSToF:
+ case BuiltinFn::kConvertUToF:
break;
}
return core::ir::Instruction::Accesses{};
diff --git a/src/tint/lang/spirv/builtin_fn.cc.tmpl b/src/tint/lang/spirv/builtin_fn.cc.tmpl
index 59c0565..8203aa9 100644
--- a/src/tint/lang/spirv/builtin_fn.cc.tmpl
+++ b/src/tint/lang/spirv/builtin_fn.cc.tmpl
@@ -105,6 +105,7 @@
case BuiltinFn::kSMod:
case BuiltinFn::kConvertFToS:
case BuiltinFn::kConvertSToF:
+ case BuiltinFn::kConvertUToF:
break;
}
return core::ir::Instruction::Accesses{};
diff --git a/src/tint/lang/spirv/builtin_fn.h b/src/tint/lang/spirv/builtin_fn.h
index 430e0a5..687eaa6 100644
--- a/src/tint/lang/spirv/builtin_fn.h
+++ b/src/tint/lang/spirv/builtin_fn.h
@@ -111,6 +111,7 @@
kSMod,
kConvertFToS,
kConvertSToF,
+ kConvertUToF,
kSdot,
kUdot,
kCooperativeMatrixLoad,
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc
index 39c200e..15f7fee 100644
--- a/src/tint/lang/spirv/intrinsic/data.cc
+++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -6711,30 +6711,37 @@
},
{
/* [63] */
+ /* fn convertUToF<R : f32_f16>[T : iu32](T) -> R */
+ /* fn convertUToF<R : f32_f16>[T : iu32, N : num](vec<N, T>) -> vec<N, R> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(180),
+ },
+ {
+ /* [64] */
/* fn sdot(u32, u32, u32) -> i32 */
/* num overloads */ 1,
/* overloads */ OverloadIndex(193),
},
{
- /* [64] */
+ /* [65] */
/* fn udot(u32, u32, u32) -> u32 */
/* num overloads */ 1,
/* overloads */ OverloadIndex(194),
},
{
- /* [65] */
+ /* [66] */
/* fn cooperative_matrix_load<T : subgroup_matrix<K, S, C, R>>[K : subgroup_matrix_kind, S : fiu32_f16, C : num, R : num](ptr<workgroup_or_storage, S, readable>, u32, u32, u32) -> T */
/* num overloads */ 1,
/* overloads */ OverloadIndex(195),
},
{
- /* [66] */
+ /* [67] */
/* fn cooperative_matrix_store[K : subgroup_matrix_kind, S : fiu32_f16, C : num, R : num](ptr<workgroup_or_storage, S, writable>, subgroup_matrix<K, S, C, R>, u32, u32, u32) */
/* num overloads */ 1,
/* overloads */ OverloadIndex(196),
},
{
- /* [67] */
+ /* [68] */
/* fn cooperative_matrix_mul_add[T : subgroup_matrix_elements, TR : subgroup_matrix_elements, C : num, R : num, K : num](subgroup_matrix<subgroup_matrix_kind_left, T, K, R>, subgroup_matrix<subgroup_matrix_kind_right, T, C, K>, subgroup_matrix<subgroup_matrix_kind_result, TR, C, R>, u32) -> subgroup_matrix<subgroup_matrix_kind_result, TR, C, R> */
/* num overloads */ 1,
/* overloads */ OverloadIndex(197),
diff --git a/src/tint/lang/spirv/reader/lower/builtins.cc b/src/tint/lang/spirv/reader/lower/builtins.cc
index 5229a87..fc2b6cc 100644
--- a/src/tint/lang/spirv/reader/lower/builtins.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins.cc
@@ -152,6 +152,9 @@
case spirv::BuiltinFn::kConvertSToF:
ConvertSToF(builtin);
break;
+ case spirv::BuiltinFn::kConvertUToF:
+ ConvertUToF(builtin);
+ break;
default:
TINT_UNREACHABLE() << "unknown spirv builtin: " << builtin->Func();
}
@@ -172,6 +175,20 @@
call->Destroy();
}
+ void ConvertUToF(spirv::ir::BuiltinCall* call) {
+ b.InsertBefore(call, [&] {
+ auto* result_ty = call->Result(0)->Type();
+
+ auto* arg = call->Args()[0];
+ if (arg->Type()->IsSignedIntegerScalarOrVector()) {
+ arg = b.Bitcast(ty.MatchWidth(ty.u32(), result_ty), arg)->Result(0);
+ }
+
+ b.ConvertWithResult(call->DetachResult(), arg);
+ });
+ call->Destroy();
+ }
+
void ConvertFToS(spirv::ir::BuiltinCall* call) {
b.InsertBefore(call, [&] {
auto* res_ty = call->Result(0)->Type();
diff --git a/src/tint/lang/spirv/reader/lower/builtins_test.cc b/src/tint/lang/spirv/reader/lower/builtins_test.cc
index 7d9de47..71ba45e 100644
--- a/src/tint/lang/spirv/reader/lower/builtins_test.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins_test.cc
@@ -5341,5 +5341,131 @@
EXPECT_EQ(expect, str());
}
+TEST_F(SpirvParser_BuiltinsTest, ConvertUToF_ScalarSigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.f32(), spirv::BuiltinFn::kConvertUToF,
+ Vector{ty.f32()}, 10_i);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:f32 = spirv.convertUToF<f32> 10i
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = bitcast 10i
+ %3:f32 = convert %2
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, ConvertUToF_ScalarUnsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.f32(), spirv::BuiltinFn::kConvertUToF,
+ Vector{ty.f32()}, 10_u);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:f32 = spirv.convertUToF<f32> 10u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:f32 = convert 10u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, ConvertUToF_VectorSigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<f32>(), spirv::BuiltinFn::kConvertUToF,
+ Vector{ty.f32()}, b.Splat<vec2<i32>>(10_i));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<f32> = spirv.convertUToF<f32> vec2<i32>(10i)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = bitcast vec2<i32>(10i)
+ %3:vec2<f32> = convert %2
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, ConvertUToF_VectorUnsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<f32>(), spirv::BuiltinFn::kConvertUToF,
+ Vector{ty.f32()}, b.Splat<vec2<u32>>(10_u));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<f32> = spirv.convertUToF<f32> vec2<u32>(10u)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<f32> = convert vec2<u32>(10u)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
} // namespace
} // namespace tint::spirv::reader::lower
diff --git a/src/tint/lang/spirv/reader/parser/convert_test.cc b/src/tint/lang/spirv/reader/parser/convert_test.cc
index ac9bcd5..c4bcacb 100644
--- a/src/tint/lang/spirv/reader/parser/convert_test.cc
+++ b/src/tint/lang/spirv/reader/parser/convert_test.cc
@@ -325,5 +325,123 @@
)");
}
+TEST_F(SpirvParserTest, ConvertUToF_ScalarSigned) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %int = OpTypeInt 32 1
+ %float = OpTypeFloat 32
+ %two = OpConstant %int 2
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpConvertUToF %float %two
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:f32 = spirv.convertUToF<f32> 2i
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, ConvertUToF_ScalarUnsigned) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %float = OpTypeFloat 32
+ %two = OpConstant %uint 2
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpConvertUToF %float %two
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:f32 = spirv.convertUToF<f32> 2u
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, ConvertUToF_VectorSigned) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %int = OpTypeInt 32 1
+ %v2int = OpTypeVector %int 2
+ %float = OpTypeFloat 32
+ %v2float = OpTypeVector %float 2
+ %two = OpConstant %int 2
+ %v2_two = OpConstantComposite %v2int %two %two
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpConvertUToF %v2float %v2_two
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<f32> = spirv.convertUToF<f32> vec2<i32>(2i)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, ConvertUToF_VectorUnsigned) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %v2uint = OpTypeVector %uint 2
+ %float = OpTypeFloat 32
+ %v2float = OpTypeVector %float 2
+ %two = OpConstant %uint 2
+ %v2_two = OpConstantComposite %v2uint %two %two
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpConvertUToF %v2float %v2_two
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<f32> = spirv.convertUToF<f32> vec2<u32>(2u)
+ ret
+ }
+}
+)");
+}
+
} // namespace
} // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index ed0c379..736ca0e 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -665,6 +665,9 @@
case spv::Op::OpConvertSToF:
EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kConvertSToF);
break;
+ case spv::Op::OpConvertUToF:
+ EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kConvertUToF);
+ break;
case spv::Op::OpAccessChain:
case spv::Op::OpInBoundsAccessChain:
EmitAccess(inst);
diff --git a/src/tint/lang/spirv/spirv.def b/src/tint/lang/spirv/spirv.def
index c74b5ce..3cf4ad5 100644
--- a/src/tint/lang/spirv/spirv.def
+++ b/src/tint/lang/spirv/spirv.def
@@ -418,6 +418,8 @@
implicit(T: f32_f16, N: num) fn convertFToS<R: iu32>(vec<N, T>) -> vec<N, R>
implicit(T: iu32) fn convertSToF<R: f32_f16>(T) -> R
implicit(T: iu32, N: num) fn convertSToF<R: f32_f16>(vec<N, T>) -> vec<N, R>
+implicit(T: iu32) fn convertUToF<R: f32_f16>(T) -> R
+implicit(T: iu32, N: num) fn convertUToF<R: f32_f16>(vec<N, T>) -> vec<N, R>
////////////////////////////////////////////////////////////////////////////////
// SPV_KHR_integer_dot_product instructions
diff --git a/src/tint/lang/spirv/writer/printer/printer.cc b/src/tint/lang/spirv/writer/printer/printer.cc
index 3d9fff1..eaed98a 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -1530,6 +1530,9 @@
case BuiltinFn::kConvertSToF:
op = spv::Op::OpConvertSToF;
break;
+ case BuiltinFn::kConvertUToF:
+ op = spv::Op::OpConvertUToF;
+ break;
case spirv::BuiltinFn::kNone:
TINT_ICE() << "undefined spirv ir function";
}
