[spirv-reader][ir] Add support for GLSL std450 FrexpStruct
Add a conversion of the FrexpStruct GLSL method to the WGSL frexp
equivalent. The spirv version allows signed or unsigned as part of the
result struct but WGSL only allows signed. Convert as needed. This is
done directly in the printer as we need to access the original SPIR-V
structure type which we'd lose otherwise.
Bug: 391673570
Change-Id: Ic9a571ed17d3a198ee052552c8f5346d43ed4665
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/224794
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/spirv/reader/import_glsl_std450_test.cc b/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
index 30b8a84..b4db5e5 100644
--- a/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
+++ b/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
@@ -77,9 +77,14 @@
%modf_result_type = OpTypeStruct %float %float
%modf_v2_result_type = OpTypeStruct %v2float %v2float
-
%ptr_function_modf_result_type = OpTypePointer Function %modf_result_type
+ %frexp_result_type_unsigned = OpTypeStruct %float %uint
+ %frexp_result_type_signed = OpTypeStruct %float %int
+ %frexp_v2_result_type_unsigned = OpTypeStruct %v2float %v2uint
+ %frexp_v2_result_type_signed = OpTypeStruct %v2float %v2int
+ %ptr_function_frexp_result_type_unsigned = OpTypePointer Function %frexp_result_type_unsigned
+
%v2uint_10_20 = OpConstantComposite %v2uint %uint_10 %uint_20
%v2uint_20_10 = OpConstantComposite %v2uint %uint_20 %uint_10
%v2uint_15_15 = OpConstantComposite %v2uint %uint_15 %uint_15
@@ -931,5 +936,173 @@
)");
}
+TEST_F(SpirvReaderTest, FrexpStruct_Store) {
+ EXPECT_IR(Preamble() + R"(
+ %1 = OpVariable %ptr_function_frexp_result_type_unsigned Function
+ %2 = OpExtInst %frexp_result_type_unsigned %glsl FrexpStruct %float_50
+ OpStore %1 %2
+ OpReturn
+ OpFunctionEnd
+ )",
+ R"(
+tint_symbol_2 = struct @align(4) {
+ tint_symbol:f32 @offset(0)
+ tint_symbol_1:u32 @offset(4)
+}
+
+__frexp_result_f32 = struct @align(4) {
+ fract:f32 @offset(0)
+ exp:i32 @offset(4)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:ptr<function, tint_symbol_2, read_write> = var
+ %3:__frexp_result_f32 = frexp 50.0f
+ %4:f32 = access %3, 0u
+ %5:i32 = access %3, 1u
+ %6:u32 = bitcast %5
+ %7:tint_symbol_2 = construct %4, %6
+ store %2, %7
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvReaderTest, FrexpStruct_ScalarUnsigned) {
+ EXPECT_IR(Preamble() + R"(
+ %1 = OpExtInst %frexp_result_type_unsigned %glsl FrexpStruct %float_50
+ %2 = OpCompositeExtract %float %1 0
+ %3 = OpCompositeExtract %uint %1 1
+ OpReturn
+ OpFunctionEnd
+ )",
+ R"(
+tint_symbol_2 = struct @align(4) {
+ tint_symbol:f32 @offset(0)
+ tint_symbol_1:u32 @offset(4)
+}
+
+__frexp_result_f32 = struct @align(4) {
+ fract:f32 @offset(0)
+ exp:i32 @offset(4)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:__frexp_result_f32 = frexp 50.0f
+ %3:f32 = access %2, 0u
+ %4:i32 = access %2, 1u
+ %5:u32 = bitcast %4
+ %6:tint_symbol_2 = construct %3, %5
+ %7:f32 = access %6, 0u
+ %8:u32 = access %6, 1u
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvReaderTest, FrexpStruct_ScalarSigned) {
+ EXPECT_IR(Preamble() + R"(
+ %1 = OpExtInst %frexp_result_type_signed %glsl FrexpStruct %float_50
+ %2 = OpCompositeExtract %float %1 0
+ %3 = OpCompositeExtract %int %1 1
+ OpReturn
+ OpFunctionEnd
+ )",
+ R"(
+tint_symbol_2 = struct @align(4) {
+ tint_symbol:f32 @offset(0)
+ tint_symbol_1:i32 @offset(4)
+}
+
+__frexp_result_f32 = struct @align(4) {
+ fract:f32 @offset(0)
+ exp:i32 @offset(4)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:__frexp_result_f32 = frexp 50.0f
+ %3:f32 = access %2, 0u
+ %4:i32 = access %2, 1u
+ %5:tint_symbol_2 = construct %3, %4
+ %6:f32 = access %5, 0u
+ %7:i32 = access %5, 1u
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvReaderTest, FrexpStruct_VectorUnsigned) {
+ EXPECT_IR(Preamble() + R"(
+ %1 = OpExtInst %frexp_v2_result_type_unsigned %glsl FrexpStruct %v2float_50_60
+ %2 = OpCompositeExtract %v2float %1 0
+ %3 = OpCompositeExtract %v2uint %1 1
+ OpReturn
+ OpFunctionEnd
+ )",
+ R"(
+tint_symbol_2 = struct @align(8) {
+ tint_symbol:vec2<f32> @offset(0)
+ tint_symbol_1:vec2<u32> @offset(8)
+}
+
+__frexp_result_vec2_f32 = struct @align(8) {
+ fract:vec2<f32> @offset(0)
+ exp:vec2<i32> @offset(8)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:__frexp_result_vec2_f32 = frexp vec2<f32>(50.0f, 60.0f)
+ %3:vec2<f32> = access %2, 0u
+ %4:vec2<i32> = access %2, 1u
+ %5:vec2<u32> = bitcast %4
+ %6:tint_symbol_2 = construct %3, %5
+ %7:vec2<f32> = access %6, 0u
+ %8:vec2<u32> = access %6, 1u
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvReaderTest, FrexpStruct_VectorSigned) {
+ EXPECT_IR(Preamble() + R"(
+ %1 = OpExtInst %frexp_v2_result_type_signed %glsl FrexpStruct %v2float_50_60
+ %2 = OpCompositeExtract %v2float %1 0
+ %3 = OpCompositeExtract %v2int %1 1
+ OpReturn
+ OpFunctionEnd
+ )",
+ R"(
+tint_symbol_2 = struct @align(8) {
+ tint_symbol:vec2<f32> @offset(0)
+ tint_symbol_1:vec2<i32> @offset(8)
+}
+
+__frexp_result_vec2_f32 = struct @align(8) {
+ fract:vec2<f32> @offset(0)
+ exp:vec2<i32> @offset(8)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:__frexp_result_vec2_f32 = frexp vec2<f32>(50.0f, 60.0f)
+ %3:vec2<f32> = access %2, 0u
+ %4:vec2<i32> = access %2, 1u
+ %5:tint_symbol_2 = construct %3, %4
+ %6:vec2<f32> = access %5, 0u
+ %7:vec2<i32> = access %5, 1u
+ 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 19eece5..149dfd6 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -704,6 +704,8 @@
return core::BuiltinFn::kClamp;
case GLSLstd450ModfStruct:
return core::BuiltinFn::kModf;
+ case GLSLstd450FrexpStruct:
+ return core::BuiltinFn::kFrexp;
case GLSLstd450NMin:
case GLSLstd450FMin: // FMin is less prescriptive about NaN operands
return core::BuiltinFn::kMin;
@@ -853,6 +855,39 @@
Emit(b_.Construct(spv_ty, fract, whole), inst.result_id());
return;
}
+ if (wgsl_fn == core::BuiltinFn::kFrexp) {
+ // For `FrexpStruct`, which is, essentially, a WGSL `frexp`
+ // instruction we need some special handling. The result type that we
+ // produce must be the SPIR-V type as we don't know how the result is
+ // used later. So, we need to make the WGSL query and re-construct an
+ // object of the right SPIR-V type. We can't, easily, do this later
+ // as we lose the SPIR-V type as soon as we replace the result of the
+ // `frexp`. So, inline the work here to generate the correct results.
+
+ auto* mem_ty = operands[0]->Type();
+ auto* result_ty = core::type::CreateFrexpResult(ty_, ir_.symbols, mem_ty);
+
+ auto* call = b_.Call(result_ty, wgsl_fn, operands);
+ auto* fract = b_.Access(mem_ty, call, 0_u);
+ auto* exp = b_.Access(ty_.MatchWidth(ty_.i32(), mem_ty), call, 1_u);
+ auto* exp_res = exp->Result(0);
+
+ EmitWithoutSpvResult(call);
+ EmitWithoutSpvResult(fract);
+ EmitWithoutSpvResult(exp);
+
+ if (auto* str = spv_ty->As<core::type::Struct>()) {
+ auto* exp_ty = str->Members()[1]->Type();
+ if (exp_ty->DeepestElement()->IsUnsignedIntegerScalar()) {
+ auto* uexp = b_.Bitcast(exp_ty, exp);
+ exp_res = uexp->Result(0);
+ EmitWithoutSpvResult(uexp);
+ }
+ }
+
+ Emit(b_.Construct(spv_ty, fract, exp_res), inst.result_id());
+ return;
+ }
if (wgsl_fn != core::BuiltinFn::kNone) {
Emit(b_.Call(spv_ty, wgsl_fn, operands), inst.result_id());
return;
