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
