Implement data unpacking intrinsics * Add support for data unpacking intrinsics * spir-v reader * type determiner * intrinsic table * spir-v, hlsl and msl writers Bug: tint:341 Change-Id: I8f40d19d59a4699af75cd579fe8398c735a77a59 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/41320 Reviewed-by: dan sinclair <dsinclair@chromium.org> Commit-Queue: Alan Baker <alanbaker@google.com>
diff --git a/src/intrinsic_table.cc b/src/intrinsic_table.cc index 0c99427..a0f3cc2 100644 --- a/src/intrinsic_table.cc +++ b/src/intrinsic_table.cc
@@ -886,127 +886,132 @@ // clang-format off - // name return type parameter types open type constraints // NOLINT - Register(I::kAbs, T, {T}, {OpenType::T, fiu32} ); // NOLINT - Register(I::kAbs, vecN_T, {vecN_T}, {OpenType::T, fiu32} ); // NOLINT - Register(I::kAcos, f32, {f32} ); // NOLINT - Register(I::kAcos, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kAll, bool_, {vecN_bool} ); // NOLINT - Register(I::kAny, bool_, {vecN_bool} ); // NOLINT - Register(I::kArrayLength, u32, {array_T} ); // NOLINT - Register(I::kAsin, f32, {f32} ); // NOLINT - Register(I::kAsin, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kAtan, f32, {f32} ); // NOLINT - Register(I::kAtan, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kAtan2, f32, {f32, f32} ); // NOLINT - Register(I::kAtan2, vecN_f32, {vecN_f32, vecN_f32} ); // NOLINT - Register(I::kCeil, f32, {f32} ); // NOLINT - Register(I::kCeil, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kClamp, T, {T, T, T}, {OpenType::T, fiu32} ); // NOLINT - Register(I::kClamp, vecN_T, {vecN_T, vecN_T, vecN_T}, {OpenType::T, fiu32} ); // NOLINT - Register(I::kCos, f32, {f32} ); // NOLINT - Register(I::kCos, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kCosh, f32, {f32} ); // NOLINT - Register(I::kCosh, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kCountOneBits, T, {T}, {OpenType::T, iu32} ); // NOLINT - Register(I::kCountOneBits, vecN_T, {vecN_T}, {OpenType::T, iu32} ); // NOLINT - Register(I::kCross, vec3_f32, {vec3_f32, vec3_f32} ); // NOLINT - Register(I::kDeterminant, f32, {matNxN_f32} ); // NOLINT - Register(I::kDistance, f32, {f32, f32} ); // NOLINT - Register(I::kDistance, f32, {vecN_f32, vecN_f32} ); // NOLINT - Register(I::kDot, f32, {vecN_f32, vecN_f32} ); // NOLINT - Register(I::kDpdx, f32, {f32} ); // NOLINT - Register(I::kDpdx, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kDpdxCoarse, f32, {f32} ); // NOLINT - Register(I::kDpdxCoarse, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kDpdxFine, f32, {f32} ); // NOLINT - Register(I::kDpdxFine, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kDpdy, f32, {f32} ); // NOLINT - Register(I::kDpdy, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kDpdyCoarse, f32, {f32} ); // NOLINT - Register(I::kDpdyCoarse, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kDpdyFine, f32, {f32} ); // NOLINT - Register(I::kDpdyFine, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kExp, f32, {f32} ); // NOLINT - Register(I::kExp, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kExp2, f32, {f32} ); // NOLINT - Register(I::kExp2, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kFaceForward, f32, {f32, f32, f32} ); // NOLINT - Register(I::kFaceForward, vecN_f32, {vecN_f32, vecN_f32, vecN_f32} ); // NOLINT - Register(I::kFloor, f32, {f32} ); // NOLINT - Register(I::kFloor, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kFma, f32, {f32, f32, f32} ); // NOLINT - Register(I::kFma, vecN_f32, {vecN_f32, vecN_f32, vecN_f32} ); // NOLINT - Register(I::kFract, f32, {f32} ); // NOLINT - Register(I::kFract, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kFrexp, f32, {f32, ptr_T}, {OpenType::T, iu32} ); // NOLINT - Register(I::kFrexp, vecN_f32, {vecN_f32, ptr_vecN_T}, {OpenType::T, iu32} ); // NOLINT - Register(I::kFwidth, f32, {f32} ); // NOLINT - Register(I::kFwidth, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kFwidthCoarse, f32, {f32} ); // NOLINT - Register(I::kFwidthCoarse, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kFwidthFine, f32, {f32} ); // NOLINT - Register(I::kFwidthFine, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kInverseSqrt, f32, {f32} ); // NOLINT - Register(I::kInverseSqrt, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kIsFinite, bool_, {f32} ); // NOLINT - Register(I::kIsFinite, vecN_bool, {vecN_f32} ); // NOLINT - Register(I::kIsInf, bool_, {f32} ); // NOLINT - Register(I::kIsInf, vecN_bool, {vecN_f32} ); // NOLINT - Register(I::kIsNan, bool_, {f32} ); // NOLINT - Register(I::kIsNan, vecN_bool, {vecN_f32} ); // NOLINT - Register(I::kIsNormal, bool_, {f32} ); // NOLINT - Register(I::kIsNormal, vecN_bool, {vecN_f32} ); // NOLINT - Register(I::kLdexp, f32, {f32, T}, {OpenType::T, iu32} ); // NOLINT - Register(I::kLdexp, vecN_f32, {vecN_f32, vecN_T}, {OpenType::T, iu32} ); // NOLINT - Register(I::kLength, f32, {f32} ); // NOLINT - Register(I::kLength, f32, {vecN_f32} ); // NOLINT - Register(I::kLog, f32, {f32} ); // NOLINT - Register(I::kLog, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kLog2, f32, {f32} ); // NOLINT - Register(I::kLog2, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kMax, T, {T, T}, {OpenType::T, fiu32} ); // NOLINT - Register(I::kMax, vecN_T, {vecN_T, vecN_T}, {OpenType::T, fiu32} ); // NOLINT - Register(I::kMin, T, {T, T}, {OpenType::T, fiu32} ); // NOLINT - Register(I::kMin, vecN_T, {vecN_T, vecN_T}, {OpenType::T, fiu32} ); // NOLINT - Register(I::kMix, f32, {f32, f32, f32} ); // NOLINT - Register(I::kMix, vecN_f32, {vecN_f32, vecN_f32, vecN_f32} ); // NOLINT - Register(I::kModf, f32, {f32, ptr_f32} ); // NOLINT - Register(I::kModf, vecN_f32, {vecN_f32, ptr_vecN_f32} ); // NOLINT - Register(I::kNormalize, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kPack2x16Float, u32, {vec2_f32} ); // NOLINT - Register(I::kPack2x16Snorm, u32, {vec2_f32} ); // NOLINT - Register(I::kPack2x16Unorm, u32, {vec2_f32} ); // NOLINT - Register(I::kPack4x8Snorm, u32, {vec4_f32} ); // NOLINT - Register(I::kPack4x8Unorm, u32, {vec4_f32} ); // NOLINT - Register(I::kPow, f32, {f32, f32} ); // NOLINT - Register(I::kPow, vecN_f32, {vecN_f32, vecN_f32} ); // NOLINT - Register(I::kReflect, f32, {f32, f32} ); // NOLINT - Register(I::kReflect, vecN_f32, {vecN_f32, vecN_f32} ); // NOLINT - Register(I::kReverseBits, T, {T}, {OpenType::T, iu32} ); // NOLINT - Register(I::kReverseBits, vecN_T, {vecN_T}, {OpenType::T, iu32} ); // NOLINT - Register(I::kRound, f32, {f32} ); // NOLINT - Register(I::kRound, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kSelect, T, {T, T, bool_}, {OpenType::T, scalar} ); // NOLINT - Register(I::kSelect, vecN_T, {vecN_T, vecN_T, vecN_bool}, {OpenType::T, scalar} ); // NOLINT - Register(I::kSign, f32, {f32} ); // NOLINT - Register(I::kSign, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kSin, f32, {f32} ); // NOLINT - Register(I::kSin, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kSinh, f32, {f32} ); // NOLINT - Register(I::kSinh, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kSmoothStep, f32, {f32, f32, f32} ); // NOLINT - Register(I::kSmoothStep, vecN_f32, {vecN_f32, vecN_f32, vecN_f32} ); // NOLINT - Register(I::kSqrt, f32, {f32} ); // NOLINT - Register(I::kSqrt, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kStep, f32, {f32, f32} ); // NOLINT - Register(I::kStep, vecN_f32, {vecN_f32, vecN_f32} ); // NOLINT - Register(I::kTan, f32, {f32} ); // NOLINT - Register(I::kTan, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kTanh, f32, {f32} ); // NOLINT - Register(I::kTanh, vecN_f32, {vecN_f32} ); // NOLINT - Register(I::kTrunc, f32, {f32} ); // NOLINT - Register(I::kTrunc, vecN_f32, {vecN_f32} ); // NOLINT + // name return type parameter types open type constraints // NOLINT + Register(I::kAbs, T, {T}, {OpenType::T, fiu32} ); // NOLINT + Register(I::kAbs, vecN_T, {vecN_T}, {OpenType::T, fiu32} ); // NOLINT + Register(I::kAcos, f32, {f32} ); // NOLINT + Register(I::kAcos, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kAll, bool_, {vecN_bool} ); // NOLINT + Register(I::kAny, bool_, {vecN_bool} ); // NOLINT + Register(I::kArrayLength, u32, {array_T} ); // NOLINT + Register(I::kAsin, f32, {f32} ); // NOLINT + Register(I::kAsin, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kAtan, f32, {f32} ); // NOLINT + Register(I::kAtan, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kAtan2, f32, {f32, f32} ); // NOLINT + Register(I::kAtan2, vecN_f32, {vecN_f32, vecN_f32} ); // NOLINT + Register(I::kCeil, f32, {f32} ); // NOLINT + Register(I::kCeil, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kClamp, T, {T, T, T}, {OpenType::T, fiu32} ); // NOLINT + Register(I::kClamp, vecN_T, {vecN_T, vecN_T, vecN_T}, {OpenType::T, fiu32} ); // NOLINT + Register(I::kCos, f32, {f32} ); // NOLINT + Register(I::kCos, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kCosh, f32, {f32} ); // NOLINT + Register(I::kCosh, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kCountOneBits, T, {T}, {OpenType::T, iu32} ); // NOLINT + Register(I::kCountOneBits, vecN_T, {vecN_T}, {OpenType::T, iu32} ); // NOLINT + Register(I::kCross, vec3_f32, {vec3_f32, vec3_f32} ); // NOLINT + Register(I::kDeterminant, f32, {matNxN_f32} ); // NOLINT + Register(I::kDistance, f32, {f32, f32} ); // NOLINT + Register(I::kDistance, f32, {vecN_f32, vecN_f32} ); // NOLINT + Register(I::kDot, f32, {vecN_f32, vecN_f32} ); // NOLINT + Register(I::kDpdx, f32, {f32} ); // NOLINT + Register(I::kDpdx, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kDpdxCoarse, f32, {f32} ); // NOLINT + Register(I::kDpdxCoarse, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kDpdxFine, f32, {f32} ); // NOLINT + Register(I::kDpdxFine, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kDpdy, f32, {f32} ); // NOLINT + Register(I::kDpdy, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kDpdyCoarse, f32, {f32} ); // NOLINT + Register(I::kDpdyCoarse, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kDpdyFine, f32, {f32} ); // NOLINT + Register(I::kDpdyFine, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kExp, f32, {f32} ); // NOLINT + Register(I::kExp, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kExp2, f32, {f32} ); // NOLINT + Register(I::kExp2, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kFaceForward, f32, {f32, f32, f32} ); // NOLINT + Register(I::kFaceForward, vecN_f32, {vecN_f32, vecN_f32, vecN_f32} ); // NOLINT + Register(I::kFloor, f32, {f32} ); // NOLINT + Register(I::kFloor, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kFma, f32, {f32, f32, f32} ); // NOLINT + Register(I::kFma, vecN_f32, {vecN_f32, vecN_f32, vecN_f32} ); // NOLINT + Register(I::kFract, f32, {f32} ); // NOLINT + Register(I::kFract, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kFrexp, f32, {f32, ptr_T}, {OpenType::T, iu32} ); // NOLINT + Register(I::kFrexp, vecN_f32, {vecN_f32, ptr_vecN_T}, {OpenType::T, iu32} ); // NOLINT + Register(I::kFwidth, f32, {f32} ); // NOLINT + Register(I::kFwidth, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kFwidthCoarse, f32, {f32} ); // NOLINT + Register(I::kFwidthCoarse, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kFwidthFine, f32, {f32} ); // NOLINT + Register(I::kFwidthFine, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kInverseSqrt, f32, {f32} ); // NOLINT + Register(I::kInverseSqrt, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kIsFinite, bool_, {f32} ); // NOLINT + Register(I::kIsFinite, vecN_bool, {vecN_f32} ); // NOLINT + Register(I::kIsInf, bool_, {f32} ); // NOLINT + Register(I::kIsInf, vecN_bool, {vecN_f32} ); // NOLINT + Register(I::kIsNan, bool_, {f32} ); // NOLINT + Register(I::kIsNan, vecN_bool, {vecN_f32} ); // NOLINT + Register(I::kIsNormal, bool_, {f32} ); // NOLINT + Register(I::kIsNormal, vecN_bool, {vecN_f32} ); // NOLINT + Register(I::kLdexp, f32, {f32, T}, {OpenType::T, iu32} ); // NOLINT + Register(I::kLdexp, vecN_f32, {vecN_f32, vecN_T}, {OpenType::T, iu32} ); // NOLINT + Register(I::kLength, f32, {f32} ); // NOLINT + Register(I::kLength, f32, {vecN_f32} ); // NOLINT + Register(I::kLog, f32, {f32} ); // NOLINT + Register(I::kLog, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kLog2, f32, {f32} ); // NOLINT + Register(I::kLog2, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kMax, T, {T, T}, {OpenType::T, fiu32} ); // NOLINT + Register(I::kMax, vecN_T, {vecN_T, vecN_T}, {OpenType::T, fiu32} ); // NOLINT + Register(I::kMin, T, {T, T}, {OpenType::T, fiu32} ); // NOLINT + Register(I::kMin, vecN_T, {vecN_T, vecN_T}, {OpenType::T, fiu32} ); // NOLINT + Register(I::kMix, f32, {f32, f32, f32} ); // NOLINT + Register(I::kMix, vecN_f32, {vecN_f32, vecN_f32, vecN_f32} ); // NOLINT + Register(I::kModf, f32, {f32, ptr_f32} ); // NOLINT + Register(I::kModf, vecN_f32, {vecN_f32, ptr_vecN_f32} ); // NOLINT + Register(I::kNormalize, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kPack2x16Float, u32, {vec2_f32} ); // NOLINT + Register(I::kPack2x16Snorm, u32, {vec2_f32} ); // NOLINT + Register(I::kPack2x16Unorm, u32, {vec2_f32} ); // NOLINT + Register(I::kPack4x8Snorm, u32, {vec4_f32} ); // NOLINT + Register(I::kPack4x8Unorm, u32, {vec4_f32} ); // NOLINT + Register(I::kPow, f32, {f32, f32} ); // NOLINT + Register(I::kPow, vecN_f32, {vecN_f32, vecN_f32} ); // NOLINT + Register(I::kReflect, f32, {f32, f32} ); // NOLINT + Register(I::kReflect, vecN_f32, {vecN_f32, vecN_f32} ); // NOLINT + Register(I::kReverseBits, T, {T}, {OpenType::T, iu32} ); // NOLINT + Register(I::kReverseBits, vecN_T, {vecN_T}, {OpenType::T, iu32} ); // NOLINT + Register(I::kRound, f32, {f32} ); // NOLINT + Register(I::kRound, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kSelect, T, {T, T, bool_}, {OpenType::T, scalar} ); // NOLINT + Register(I::kSelect, vecN_T, {vecN_T, vecN_T, vecN_bool}, {OpenType::T, scalar} ); // NOLINT + Register(I::kSign, f32, {f32} ); // NOLINT + Register(I::kSign, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kSin, f32, {f32} ); // NOLINT + Register(I::kSin, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kSinh, f32, {f32} ); // NOLINT + Register(I::kSinh, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kSmoothStep, f32, {f32, f32, f32} ); // NOLINT + Register(I::kSmoothStep, vecN_f32, {vecN_f32, vecN_f32, vecN_f32} ); // NOLINT + Register(I::kSqrt, f32, {f32} ); // NOLINT + Register(I::kSqrt, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kStep, f32, {f32, f32} ); // NOLINT + Register(I::kStep, vecN_f32, {vecN_f32, vecN_f32} ); // NOLINT + Register(I::kTan, f32, {f32} ); // NOLINT + Register(I::kTan, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kTanh, f32, {f32} ); // NOLINT + Register(I::kTanh, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kTrunc, f32, {f32} ); // NOLINT + Register(I::kTrunc, vecN_f32, {vecN_f32} ); // NOLINT + Register(I::kUnpack2x16Float, vec2_f32, {u32} ); // NOLINT + Register(I::kUnpack2x16Snorm, vec2_f32, {u32} ); // NOLINT + Register(I::kUnpack2x16Unorm, vec2_f32, {u32} ); // NOLINT + Register(I::kUnpack4x8Snorm, vec4_f32, {u32} ); // NOLINT + Register(I::kUnpack4x8Unorm, vec4_f32, {u32} ); // NOLINT // clang-format on auto* tex_1d_f32 = sampled_texture(Dim::k1d, f32);
diff --git a/src/reader/spirv/function.cc b/src/reader/spirv/function.cc index 46fb806..b68b9c9 100644 --- a/src/reader/spirv/function.cc +++ b/src/reader/spirv/function.cc
@@ -425,6 +425,16 @@ return "tanh"; case GLSLstd450Trunc: return "trunc"; + case GLSLstd450UnpackSnorm4x8: + return "unpack4x8snorm"; + case GLSLstd450UnpackUnorm4x8: + return "unpack4x8unorm"; + case GLSLstd450UnpackSnorm2x16: + return "unpack2x16snorm"; + case GLSLstd450UnpackUnorm2x16: + return "unpack2x16unorm"; + case GLSLstd450UnpackHalf2x16: + return "unpack2x16float"; default: // TODO(dneto) - The following are not implemented. @@ -448,11 +458,6 @@ case GLSLstd450FrexpStruct: case GLSLstd450PackDouble2x32: - case GLSLstd450UnpackSnorm2x16: - case GLSLstd450UnpackUnorm2x16: - case GLSLstd450UnpackHalf2x16: - case GLSLstd450UnpackSnorm4x8: - case GLSLstd450UnpackUnorm4x8: case GLSLstd450UnpackDouble2x32: case GLSLstd450Refract:
diff --git a/src/reader/spirv/function_glsl_std_450_test.cc b/src/reader/spirv/function_glsl_std_450_test.cc index b876e9a..d0c43f1 100644 --- a/src/reader/spirv/function_glsl_std_450_test.cc +++ b/src/reader/spirv/function_glsl_std_450_test.cc
@@ -1548,6 +1548,50 @@ {"PackUnorm2x16", "pack2x16unorm", 2}, {"PackHalf2x16", "pack2x16float", 2}})); +using SpvParserTest_GlslStd450_DataUnpacking = + SpvParserTestBase<::testing::TestWithParam<DataPackingCase>>; + +TEST_P(SpvParserTest_GlslStd450_DataUnpacking, Valid) { + auto param = GetParam(); + const auto assembly = Preamble() + R"( + %1 = OpExtInst )" + (param.vec_size == 2 ? "%v2float" : "%v4float") + + std::string(" %glsl ") + param.opcode + R"( %u1 + OpReturn + OpFunctionEnd + )"; + auto p = parser(test::Assemble(assembly)); + ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << assembly; + FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100)); + EXPECT_TRUE(fe.EmitBody()) << p->error(); + const auto body = ToString(p->builder(), fe.ast_body()); + EXPECT_THAT(body, HasSubstr(R"( + VariableConst{ + x_1 + none + )" + std::string(param.vec_size == 2 ? "__vec_2__f32" : "__vec_4__f32") + + R"( + { + Call[not set]{ + Identifier[not set]{)" + + param.wgsl_func + R"(} + ( + Identifier[not set]{u1} + ) + } + } + })")) + << body; +} + +INSTANTIATE_TEST_SUITE_P(Samples, + SpvParserTest_GlslStd450_DataUnpacking, + ::testing::ValuesIn(std::vector<DataPackingCase>{ + {"UnpackSnorm4x8", "unpack4x8snorm", 4}, + {"UnpackUnorm4x8", "unpack4x8unorm", 4}, + {"UnpackSnorm2x16", "unpack2x16snorm", 2}, + {"UnpackUnorm2x16", "unpack2x16unorm", 2}, + {"UnpackHalf2x16", "unpack2x16float", 2}})); + } // namespace } // namespace spirv } // namespace reader
diff --git a/src/semantic/intrinsic.h b/src/semantic/intrinsic.h index 6cc116b..d8a313b 100644 --- a/src/semantic/intrinsic.h +++ b/src/semantic/intrinsic.h
@@ -101,7 +101,12 @@ kTextureSampleGrad, kTextureSampleLevel, kTextureStore, - kTrunc + kTrunc, + kUnpack4x8Snorm, + kUnpack4x8Unorm, + kUnpack2x16Snorm, + kUnpack2x16Unorm, + kUnpack2x16Float, }; /// @returns the name of the intrinsic function type. The spelling, including @@ -143,6 +148,11 @@ /// @returns true if the given `i` is a data packing intrinsic bool IsDataPackingIntrinsic(IntrinsicType i); +/// Determines if the given `i` is a data unpacking intrinsic +/// @param i the intrinsic +/// @returns true if the given `i` is a data unpacking intrinsic +bool IsDataUnpackingIntrinsic(IntrinsicType i); + /// Intrinsic holds the semantic information for an intrinsic function. class Intrinsic : public Castable<Intrinsic, CallTarget> { public: @@ -185,6 +195,9 @@ /// @returns true if intrinsic is a data packing intrinsic bool IsDataPacking() const; + /// @returns true if intrinsic is a data unpacking intrinsic + bool IsDataUnpacking() const; + private: IntrinsicType const type_; };
diff --git a/src/semantic/sem_intrinsic.cc b/src/semantic/sem_intrinsic.cc index aa398c0..9e610ee 100644 --- a/src/semantic/sem_intrinsic.cc +++ b/src/semantic/sem_intrinsic.cc
@@ -188,6 +188,16 @@ return "textureStore"; case IntrinsicType::kTrunc: return "trunc"; + case IntrinsicType::kUnpack4x8Snorm: + return "unpack4x8snorm"; + case IntrinsicType::kUnpack4x8Unorm: + return "unpack4x8unorm"; + case IntrinsicType::kUnpack2x16Snorm: + return "unpack2x16snorm"; + case IntrinsicType::kUnpack2x16Unorm: + return "unpack2x16unorm"; + case IntrinsicType::kUnpack2x16Float: + return "unpack2x16float"; } return "<unknown>"; } @@ -238,6 +248,14 @@ i == IntrinsicType::kPack2x16Float; } +bool IsDataUnpackingIntrinsic(IntrinsicType i) { + return i == IntrinsicType::kUnpack4x8Snorm || + i == IntrinsicType::kUnpack4x8Unorm || + i == IntrinsicType::kUnpack2x16Snorm || + i == IntrinsicType::kUnpack2x16Unorm || + i == IntrinsicType::kUnpack2x16Float; +} + Intrinsic::Intrinsic(IntrinsicType type, type::Type* return_type, const ParameterList& parameters) @@ -273,5 +291,9 @@ return IsDataPackingIntrinsic(type_); } +bool Intrinsic::IsDataUnpacking() const { + return IsDataUnpackingIntrinsic(type_); +} + } // namespace semantic } // namespace tint
diff --git a/src/type_determiner.cc b/src/type_determiner.cc index 18f52f6..a127caa 100644 --- a/src/type_determiner.cc +++ b/src/type_determiner.cc
@@ -717,6 +717,16 @@ return IntrinsicType::kTextureSampleLevel; } else if (name == "trunc") { return IntrinsicType::kTrunc; + } else if (name == "unpack4x8snorm") { + return IntrinsicType::kUnpack4x8Snorm; + } else if (name == "unpack4x8unorm") { + return IntrinsicType::kUnpack4x8Unorm; + } else if (name == "unpack2x16snorm") { + return IntrinsicType::kUnpack2x16Snorm; + } else if (name == "unpack2x16unorm") { + return IntrinsicType::kUnpack2x16Unorm; + } else if (name == "unpack2x16float") { + return IntrinsicType::kUnpack2x16Float; } return IntrinsicType::kNone; }
diff --git a/src/type_determiner_test.cc b/src/type_determiner_test.cc index 71e0371..fc46468 100644 --- a/src/type_determiner_test.cc +++ b/src/type_determiner_test.cc
@@ -1794,6 +1794,36 @@ IntrinsicData{"pack2x16unorm", IntrinsicType::kPack2x16Unorm}, IntrinsicData{"pack2x16float", IntrinsicType::kPack2x16Float})); +using ImportData_DataUnpackingTest = TypeDeterminerTestWithParam<IntrinsicData>; +TEST_P(ImportData_DataUnpackingTest, InferType) { + auto param = GetParam(); + + bool pack4 = param.intrinsic == IntrinsicType::kUnpack4x8Snorm || + param.intrinsic == IntrinsicType::kUnpack4x8Unorm; + + auto* call = Call(param.name, 1u); + WrapInFunction(call); + + EXPECT_TRUE(td()->Determine()) << td()->error(); + ASSERT_NE(TypeOf(call), nullptr); + EXPECT_TRUE(TypeOf(call)->is_float_vector()); + if (pack4) { + EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 4u); + } else { + EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 2u); + } +} + +INSTANTIATE_TEST_SUITE_P( + TypeDeterminerTest, + ImportData_DataUnpackingTest, + testing::Values( + IntrinsicData{"unpack4x8snorm", IntrinsicType::kUnpack4x8Snorm}, + IntrinsicData{"unpack4x8unorm", IntrinsicType::kUnpack4x8Unorm}, + IntrinsicData{"unpack2x16snorm", IntrinsicType::kUnpack2x16Snorm}, + IntrinsicData{"unpack2x16unorm", IntrinsicType::kUnpack2x16Unorm}, + IntrinsicData{"unpack2x16float", IntrinsicType::kUnpack2x16Float})); + using ImportData_SingleParamTest = TypeDeterminerTestWithParam<IntrinsicData>; TEST_P(ImportData_SingleParamTest, Scalar) { auto param = GetParam();
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc index d370a62..b170fbd 100644 --- a/src/writer/hlsl/generator_impl.cc +++ b/src/writer/hlsl/generator_impl.cc
@@ -556,6 +556,8 @@ return false; } else if (intrinsic->IsDataPacking()) { return EmitDataPackingCall(pre, out, expr, intrinsic); + } else if (intrinsic->IsDataUnpacking()) { + return EmitDataUnpackingCall(pre, out, expr, intrinsic); } auto name = generate_builtin_name(intrinsic); if (name.empty()) { @@ -694,6 +696,77 @@ return true; } +bool GeneratorImpl::EmitDataUnpackingCall( + std::ostream& pre, + std::ostream& out, + ast::CallExpression* expr, + const semantic::Intrinsic* intrinsic) { + auto* param = expr->params()[0]; + auto tmp_name = generate_name(kTempNamePrefix); + std::ostringstream expr_out; + if (!EmitExpression(pre, expr_out, param)) { + return false; + } + uint32_t dims = 2; + bool is_signed = false; + uint32_t scale = 65535; + if (intrinsic->Type() == semantic::IntrinsicType::kUnpack4x8Snorm || + intrinsic->Type() == semantic::IntrinsicType::kUnpack4x8Unorm) { + dims = 4; + scale = 255; + } + if (intrinsic->Type() == semantic::IntrinsicType::kUnpack4x8Snorm || + intrinsic->Type() == semantic::IntrinsicType::kUnpack2x16Snorm) { + is_signed = true; + scale = (scale - 1) / 2; + } + switch (intrinsic->Type()) { + case semantic::IntrinsicType::kUnpack4x8Snorm: + case semantic::IntrinsicType::kUnpack2x16Snorm: { + auto tmp_name2 = generate_name(kTempNamePrefix); + pre << "int " << tmp_name2 << " = int(" << expr_out.str() << ");\n"; + // Perform sign extension on the converted values. + pre << "int" << dims << " " << tmp_name << " = int" << dims << "("; + if (dims == 2) { + pre << tmp_name2 << " << 16, " << tmp_name2 << ") >> 16"; + } else { + pre << tmp_name2 << " << 24, " << tmp_name2 << " << 16, " << tmp_name2 + << " << 8, " << tmp_name2 << ") >> 24"; + } + pre << ";\n"; + out << "clamp(float" << dims << "(" << tmp_name << ") / " << scale + << ".0, " << (is_signed ? "-1.0" : "0.0") << ", 1.0)"; + break; + } + case semantic::IntrinsicType::kUnpack4x8Unorm: + case semantic::IntrinsicType::kUnpack2x16Unorm: { + auto tmp_name2 = generate_name(kTempNamePrefix); + pre << "uint " << tmp_name2 << " = " << expr_out.str() << ";\n"; + pre << "uint" << dims << " " << tmp_name << " = uint" << dims << "("; + pre << tmp_name2 << " & " << (dims == 2 ? "0xffff" : "0xff") << ", "; + if (dims == 4) { + pre << "(" << tmp_name2 << " >> " << (32 / dims) << ") & 0xff, (" + << tmp_name2 << " >> 16) & 0xff, " << tmp_name2 << " >> 24"; + } else { + pre << tmp_name2 << " >> " << (32 / dims); + } + pre << ");\n"; + out << "float" << dims << "(" << tmp_name << ") / " << scale << ".0"; + break; + } + case semantic::IntrinsicType::kUnpack2x16Float: + pre << "uint " << tmp_name << " = " << expr_out.str() << ";\n"; + out << "f16tof32(uint2(" << tmp_name << " & 0xffff, " << tmp_name + << " >> 16))"; + break; + default: + error_ = "Internal error: unhandled data packing intrinsic"; + return false; + } + + return true; +} + bool GeneratorImpl::EmitTextureCall(std::ostream& pre, std::ostream& out, ast::CallExpression* expr,
diff --git a/src/writer/hlsl/generator_impl.h b/src/writer/hlsl/generator_impl.h index 805a0ce..23ac598 100644 --- a/src/writer/hlsl/generator_impl.h +++ b/src/writer/hlsl/generator_impl.h
@@ -169,6 +169,16 @@ std::ostream& out, ast::CallExpression* expr, const semantic::Intrinsic* intrinsic); + /// Handles generating a call to data unpacking intrinsic + /// @param pre the preamble of the expression stream + /// @param out the output of the expression stream + /// @param expr the call expression + /// @param intrinsic the semantic information for the texture intrinsic + /// @returns true if the call expression is emitted + bool EmitDataUnpackingCall(std::ostream& pre, + std::ostream& out, + ast::CallExpression* expr, + const semantic::Intrinsic* intrinsic); /// Handles a case statement /// @param out the output stream /// @param stmt the statement
diff --git a/src/writer/hlsl/generator_impl_intrinsic_test.cc b/src/writer/hlsl/generator_impl_intrinsic_test.cc index 2b32853..4ec2683 100644 --- a/src/writer/hlsl/generator_impl_intrinsic_test.cc +++ b/src/writer/hlsl/generator_impl_intrinsic_test.cc
@@ -330,7 +330,7 @@ EXPECT_THAT(result(), HasSubstr("(_tint_tmp.x | _tint_tmp.y << 16)")); } -TEST_F(HlslGeneratorImplTest_Intrinsic, Pack2x16float) { +TEST_F(HlslGeneratorImplTest_Intrinsic, Pack2x16Float) { auto* call = Call("pack2x16float", "p1"); Global("p1", ast::StorageClass::kPrivate, ty.vec2<f32>()); WrapInFunction(call); @@ -342,6 +342,85 @@ EXPECT_THAT(result(), HasSubstr("(_tint_tmp.x | _tint_tmp.y << 16)")); } +TEST_F(HlslGeneratorImplTest_Intrinsic, Unpack4x8Snorm) { + auto* call = Call("unpack4x8snorm", "p1"); + Global("p1", ast::StorageClass::kPrivate, ty.u32()); + WrapInFunction(call); + GeneratorImpl& gen = Build(); + + gen.increment_indent(); + ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error(); + EXPECT_THAT(pre_result(), HasSubstr("int _tint_tmp_0 = int(p1);")); + EXPECT_THAT(pre_result(), + HasSubstr("int4 _tint_tmp = int4(_tint_tmp_0 << 24, _tint_tmp_0 " + "<< 16, _tint_tmp_0 << 8, _tint_tmp_0) >> 24;")); + EXPECT_THAT(result(), + HasSubstr("clamp(float4(_tint_tmp) / 127.0, -1.0, 1.0)")); +} + +TEST_F(HlslGeneratorImplTest_Intrinsic, Unpack4x8Unorm) { + auto* call = Call("unpack4x8unorm", "p1"); + Global("p1", ast::StorageClass::kPrivate, ty.u32()); + WrapInFunction(call); + GeneratorImpl& gen = Build(); + + gen.increment_indent(); + ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error(); + EXPECT_THAT(pre_result(), HasSubstr("uint _tint_tmp_0 = p1;")); + EXPECT_THAT( + pre_result(), + HasSubstr("uint4 _tint_tmp = uint4(_tint_tmp_0 & 0xff, (_tint_tmp_0 >> " + "8) & 0xff, (_tint_tmp_0 >> 16) & 0xff, _tint_tmp_0 >> 24);")); + EXPECT_THAT(result(), HasSubstr("float4(_tint_tmp) / 255.0")); +} + +TEST_F(HlslGeneratorImplTest_Intrinsic, Unpack2x16Snorm) { + auto* call = Call("unpack2x16snorm", "p1"); + Global("p1", ast::StorageClass::kPrivate, ty.u32()); + WrapInFunction(call); + GeneratorImpl& gen = Build(); + + gen.increment_indent(); + ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error(); + EXPECT_THAT(pre_result(), HasSubstr("int _tint_tmp_0 = int(p1);")); + EXPECT_THAT( + pre_result(), + HasSubstr( + "int2 _tint_tmp = int2(_tint_tmp_0 << 16, _tint_tmp_0) >> 16;")); + EXPECT_THAT(result(), + HasSubstr("clamp(float2(_tint_tmp) / 32767.0, -1.0, 1.0)")); +} + +TEST_F(HlslGeneratorImplTest_Intrinsic, Unpack2x16Unorm) { + auto* call = Call("unpack2x16unorm", "p1"); + Global("p1", ast::StorageClass::kPrivate, ty.u32()); + WrapInFunction(call); + GeneratorImpl& gen = Build(); + + gen.increment_indent(); + ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error(); + EXPECT_THAT(pre_result(), HasSubstr("uint _tint_tmp_0 = p1;")); + EXPECT_THAT( + pre_result(), + HasSubstr( + "uint2 _tint_tmp = uint2(_tint_tmp_0 & 0xffff, _tint_tmp_0 >> 16);")); + EXPECT_THAT(result(), HasSubstr("float2(_tint_tmp) / 65535.0")); +} + +TEST_F(HlslGeneratorImplTest_Intrinsic, Unpack2x16Float) { + auto* call = Call("unpack2x16float", "p1"); + Global("p1", ast::StorageClass::kPrivate, ty.u32()); + WrapInFunction(call); + GeneratorImpl& gen = Build(); + + gen.increment_indent(); + ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error(); + EXPECT_THAT(pre_result(), HasSubstr("uint _tint_tmp = p1;")); + EXPECT_THAT( + result(), + HasSubstr("f16tof32(uint2(_tint_tmp & 0xffff, _tint_tmp >> 16))")); +} + } // namespace } // namespace hlsl } // namespace writer
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc index 1f71561..5421a64 100644 --- a/src/writer/msl/generator_impl.cc +++ b/src/writer/msl/generator_impl.cc
@@ -449,9 +449,14 @@ if (intrinsic->IsTexture()) { return EmitTextureCall(expr, intrinsic); } - if (intrinsic->Type() == semantic::IntrinsicType::kPack2x16Float) { + if (intrinsic->Type() == semantic::IntrinsicType::kPack2x16Float || + intrinsic->Type() == semantic::IntrinsicType::kUnpack2x16Float) { make_indent(); - out_ << "as_type<uint>(half2("; + if (intrinsic->Type() == semantic::IntrinsicType::kPack2x16Float) { + out_ << "as_type<uint>(half2("; + } else { + out_ << "float2(as_type<half2>("; + } if (!EmitExpression(expr->params()[0])) { return false; } @@ -900,6 +905,18 @@ case semantic::IntrinsicType::kInverseSqrt: out += "rsqrt"; break; + case semantic::IntrinsicType::kUnpack4x8Snorm: + out += "unpack_snorm4x8_to_float"; + break; + case semantic::IntrinsicType::kUnpack4x8Unorm: + out += "unpack_unorm4x8_to_float"; + break; + case semantic::IntrinsicType::kUnpack2x16Snorm: + out += "unpack_snorm2x16_to_float"; + break; + case semantic::IntrinsicType::kUnpack2x16Unorm: + out += "unpack_unorm2x16_to_float"; + break; default: error_ = "Unknown import method: " + std::string(intrinsic->str()); return "";
diff --git a/src/writer/msl/generator_impl_intrinsic_test.cc b/src/writer/msl/generator_impl_intrinsic_test.cc index 56b194a..3873b4c 100644 --- a/src/writer/msl/generator_impl_intrinsic_test.cc +++ b/src/writer/msl/generator_impl_intrinsic_test.cc
@@ -157,6 +157,11 @@ case IntrinsicType::kPack4x8Snorm: case IntrinsicType::kPack4x8Unorm: return builder->Call(str.str(), "f4"); + case IntrinsicType::kUnpack4x8Snorm: + case IntrinsicType::kUnpack4x8Unorm: + case IntrinsicType::kUnpack2x16Snorm: + case IntrinsicType::kUnpack2x16Unorm: + return builder->Call(str.str(), "u1"); default: break; } @@ -174,6 +179,7 @@ Global("f2", ast::StorageClass::kFunction, ty.vec2<float>()); Global("f3", ast::StorageClass::kFunction, ty.vec3<float>()); Global("f4", ast::StorageClass::kFunction, ty.vec4<float>()); + Global("u1", ast::StorageClass::kFunction, ty.u32()); Global("u2", ast::StorageClass::kFunction, ty.vec2<unsigned int>()); Global("b2", ast::StorageClass::kFunction, ty.vec2<bool>()); Global("m2x2", ast::StorageClass::kFunction, ty.mat2x2<float>()); @@ -276,7 +282,15 @@ IntrinsicData{IntrinsicType::kStep, ParamType::kF32, "metal::step"}, IntrinsicData{IntrinsicType::kTan, ParamType::kF32, "metal::tan"}, IntrinsicData{IntrinsicType::kTanh, ParamType::kF32, "metal::tanh"}, - IntrinsicData{IntrinsicType::kTrunc, ParamType::kF32, "metal::trunc"})); + IntrinsicData{IntrinsicType::kTrunc, ParamType::kF32, "metal::trunc"}, + IntrinsicData{IntrinsicType::kUnpack4x8Snorm, ParamType::kU32, + "metal::unpack_snorm4x8_to_float"}, + IntrinsicData{IntrinsicType::kUnpack4x8Unorm, ParamType::kU32, + "metal::unpack_unorm4x8_to_float"}, + IntrinsicData{IntrinsicType::kUnpack2x16Snorm, ParamType::kU32, + "metal::unpack_snorm2x16_to_float"}, + IntrinsicData{IntrinsicType::kUnpack2x16Unorm, ParamType::kU32, + "metal::unpack_unorm2x16_to_float"})); TEST_F(MslGeneratorImplTest, Intrinsic_Call) { Global("param1", ast::StorageClass::kFunction, ty.vec2<f32>()); @@ -304,6 +318,18 @@ EXPECT_EQ(gen.result(), " as_type<uint>(half2(p1))"); } +TEST_F(MslGeneratorImplTest, Unpack2x16Float) { + auto* call = Call("unpack2x16float", "p1"); + Global("p1", ast::StorageClass::kFunction, ty.u32()); + WrapInFunction(call); + + GeneratorImpl& gen = Build(); + + gen.increment_indent(); + ASSERT_TRUE(gen.EmitExpression(call)) << gen.error(); + EXPECT_EQ(gen.result(), " float2(as_type<half2>(p1))"); +} + } // namespace } // namespace msl } // namespace writer
diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc index 6bdeab5..de8a747 100644 --- a/src/writer/spirv/builder.cc +++ b/src/writer/spirv/builder.cc
@@ -284,6 +284,16 @@ return GLSLstd450Tanh; case IntrinsicType::kTrunc: return GLSLstd450Trunc; + case IntrinsicType::kUnpack4x8Snorm: + return GLSLstd450UnpackSnorm4x8; + case IntrinsicType::kUnpack4x8Unorm: + return GLSLstd450UnpackUnorm4x8; + case IntrinsicType::kUnpack2x16Snorm: + return GLSLstd450UnpackSnorm2x16; + case IntrinsicType::kUnpack2x16Unorm: + return GLSLstd450UnpackUnorm2x16; + case IntrinsicType::kUnpack2x16Float: + return GLSLstd450UnpackHalf2x16; default: break; }
diff --git a/src/writer/spirv/builder_intrinsic_test.cc b/src/writer/spirv/builder_intrinsic_test.cc index 1095e34..b20e553 100644 --- a/src/writer/spirv/builder_intrinsic_test.cc +++ b/src/writer/spirv/builder_intrinsic_test.cc
@@ -1506,6 +1506,131 @@ Validate(b); } +using Intrinsic_Builtin_DataPacking_Test = + IntrinsicBuilderTestWithParam<IntrinsicData>; +TEST_P(Intrinsic_Builtin_DataPacking_Test, Binary) { + auto param = GetParam(); + + bool pack4 = param.name == "pack4x8snorm" || param.name == "pack4x8unorm"; + auto* call = pack4 ? Call(param.name, vec4<float>(1.0f, 1.0f, 1.0f, 1.0f)) + : Call(param.name, vec2<float>(1.0f, 1.0f)); + WrapInFunction(call); + + auto* func = Func("a_func", ast::VariableList{}, ty.void_(), + ast::StatementList{}, ast::FunctionDecorationList{}); + + spirv::Builder& b = Build(); + + ASSERT_TRUE(b.GenerateFunction(func)) << b.error(); + + EXPECT_EQ(b.GenerateCallExpression(call), 5u) << b.error(); + if (pack4) { + EXPECT_EQ(DumpBuilder(b), R"(%7 = OpExtInstImport "GLSL.std.450" +OpName %3 "a_func" +%2 = OpTypeVoid +%1 = OpTypeFunction %2 +%6 = OpTypeInt 32 0 +%9 = OpTypeFloat 32 +%8 = OpTypeVector %9 4 +%10 = OpConstant %9 1 +%11 = OpConstantComposite %8 %10 %10 %10 %10 +%3 = OpFunction %2 None %1 +%4 = OpLabel +%5 = OpExtInst %6 %7 )" + param.op + + R"( %11 +OpReturn +OpFunctionEnd +)"); + } else { + EXPECT_EQ(DumpBuilder(b), R"(%7 = OpExtInstImport "GLSL.std.450" +OpName %3 "a_func" +%2 = OpTypeVoid +%1 = OpTypeFunction %2 +%6 = OpTypeInt 32 0 +%9 = OpTypeFloat 32 +%8 = OpTypeVector %9 2 +%10 = OpConstant %9 1 +%11 = OpConstantComposite %8 %10 %10 +%3 = OpFunction %2 None %1 +%4 = OpLabel +%5 = OpExtInst %6 %7 )" + param.op + + R"( %11 +OpReturn +OpFunctionEnd +)"); + } +} + +INSTANTIATE_TEST_SUITE_P( + IntrinsicBuilderTest, + Intrinsic_Builtin_DataPacking_Test, + testing::Values(IntrinsicData{"pack4x8snorm", "PackSnorm4x8"}, + IntrinsicData{"pack4x8unorm", "PackUnorm4x8"}, + IntrinsicData{"pack2x16snorm", "PackSnorm2x16"}, + IntrinsicData{"pack2x16unorm", "PackUnorm2x16"}, + IntrinsicData{"pack2x16float", "PackHalf2x16"})); + +using Intrinsic_Builtin_DataUnpacking_Test = + IntrinsicBuilderTestWithParam<IntrinsicData>; +TEST_P(Intrinsic_Builtin_DataUnpacking_Test, Binary) { + auto param = GetParam(); + + bool pack4 = param.name == "unpack4x8snorm" || param.name == "unpack4x8unorm"; + auto* call = Call(param.name, 1u); + WrapInFunction(call); + + auto* func = Func("a_func", ast::VariableList{}, ty.void_(), + ast::StatementList{}, ast::FunctionDecorationList{}); + + spirv::Builder& b = Build(); + + ASSERT_TRUE(b.GenerateFunction(func)) << b.error(); + + EXPECT_EQ(b.GenerateCallExpression(call), 5u) << b.error(); + if (pack4) { + EXPECT_EQ(DumpBuilder(b), R"(%8 = OpExtInstImport "GLSL.std.450" +OpName %3 "a_func" +%2 = OpTypeVoid +%1 = OpTypeFunction %2 +%7 = OpTypeFloat 32 +%6 = OpTypeVector %7 4 +%9 = OpTypeInt 32 0 +%10 = OpConstant %9 1 +%3 = OpFunction %2 None %1 +%4 = OpLabel +%5 = OpExtInst %6 %8 )" + param.op + + R"( %10 +OpReturn +OpFunctionEnd +)"); + } else { + EXPECT_EQ(DumpBuilder(b), R"(%8 = OpExtInstImport "GLSL.std.450" +OpName %3 "a_func" +%2 = OpTypeVoid +%1 = OpTypeFunction %2 +%7 = OpTypeFloat 32 +%6 = OpTypeVector %7 2 +%9 = OpTypeInt 32 0 +%10 = OpConstant %9 1 +%3 = OpFunction %2 None %1 +%4 = OpLabel +%5 = OpExtInst %6 %8 )" + param.op + + R"( %10 +OpReturn +OpFunctionEnd +)"); + } +} + +INSTANTIATE_TEST_SUITE_P( + IntrinsicBuilderTest, + Intrinsic_Builtin_DataUnpacking_Test, + testing::Values(IntrinsicData{"unpack4x8snorm", "UnpackSnorm4x8"}, + IntrinsicData{"unpack4x8unorm", "UnpackUnorm4x8"}, + IntrinsicData{"unpack2x16snorm", "UnpackSnorm2x16"}, + IntrinsicData{"unpack2x16unorm", "UnpackUnorm2x16"}, + IntrinsicData{"unpack2x16float", "UnpackHalf2x16"})); + } // namespace } // namespace spirv } // namespace writer