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>
This commit is contained in:
Alan Baker
Commit Bot service account
parent
2a284b2a13
commit
cd17ea88e3
|
|
@ -886,127 +886,132 @@ Impl::Impl() {
|
|||
|
||||
// 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);
|
||||
|
|
|
|||
|
|
@ -425,6 +425,16 @@ std::string GetGlslStd450FuncName(uint32_t ext_opcode) {
|
|||
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 @@ std::string GetGlslStd450FuncName(uint32_t ext_opcode) {
|
|||
case GLSLstd450FrexpStruct:
|
||||
|
||||
case GLSLstd450PackDouble2x32:
|
||||
case GLSLstd450UnpackSnorm2x16:
|
||||
case GLSLstd450UnpackUnorm2x16:
|
||||
case GLSLstd450UnpackHalf2x16:
|
||||
case GLSLstd450UnpackSnorm4x8:
|
||||
case GLSLstd450UnpackUnorm4x8:
|
||||
case GLSLstd450UnpackDouble2x32:
|
||||
|
||||
case GLSLstd450Refract:
|
||||
|
|
|
|||
|
|
@ -1548,6 +1548,50 @@ INSTANTIATE_TEST_SUITE_P(Samples,
|
|||
{"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
|
||||
|
|
|
|||
|
|
@ -101,7 +101,12 @@ enum class IntrinsicType {
|
|||
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 @@ bool IsImageQueryIntrinsic(IntrinsicType i);
|
|||
/// @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 @@ class Intrinsic : public Castable<Intrinsic, CallTarget> {
|
|||
/// @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_;
|
||||
};
|
||||
|
|
|
|||
|
|
@ -188,6 +188,16 @@ const char* str(IntrinsicType i) {
|
|||
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 @@ bool IsDataPackingIntrinsic(IntrinsicType i) {
|
|||
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 @@ bool Intrinsic::IsDataPacking() const {
|
|||
return IsDataPackingIntrinsic(type_);
|
||||
}
|
||||
|
||||
bool Intrinsic::IsDataUnpacking() const {
|
||||
return IsDataUnpackingIntrinsic(type_);
|
||||
}
|
||||
|
||||
} // namespace semantic
|
||||
} // namespace tint
|
||||
|
|
|
|||
|
|
@ -717,6 +717,16 @@ IntrinsicType TypeDeterminer::MatchIntrinsicType(const std::string& name) {
|
|||
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;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -1794,6 +1794,36 @@ INSTANTIATE_TEST_SUITE_P(
|
|||
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();
|
||||
|
|
|
|||
|
|
@ -556,6 +556,8 @@ bool GeneratorImpl::EmitCall(std::ostream& pre,
|
|||
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 @@ bool GeneratorImpl::EmitDataPackingCall(std::ostream& pre,
|
|||
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,
|
||||
|
|
|
|||
|
|
@ -169,6 +169,16 @@ class GeneratorImpl {
|
|||
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
|
||||
|
|
|
|||
|
|
@ -330,7 +330,7 @@ TEST_F(HlslGeneratorImplTest_Intrinsic, Pack2x16Unorm) {
|
|||
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 @@ TEST_F(HlslGeneratorImplTest_Intrinsic, Pack2x16float) {
|
|||
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
|
||||
|
|
|
|||
|
|
@ -449,9 +449,14 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
|
|||
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 @@ std::string GeneratorImpl::generate_builtin_name(
|
|||
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 "";
|
||||
|
|
|
|||
|
|
@ -157,6 +157,11 @@ ast::CallExpression* GenerateCall(IntrinsicType intrinsic,
|
|||
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 @@ TEST_P(MslIntrinsicTest, Emit) {
|
|||
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 @@ INSTANTIATE_TEST_SUITE_P(
|
|||
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 @@ TEST_F(MslGeneratorImplTest, Pack2x16Float) {
|
|||
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
|
||||
|
|
|
|||
|
|
@ -284,6 +284,16 @@ uint32_t intrinsic_to_glsl_method(const semantic::Intrinsic* intrinsic) {
|
|||
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;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -1506,6 +1506,131 @@ TEST_F(IntrinsicBuilderTest, Call_ArrayLength_OtherMembersInStruct) {
|
|||
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
|
||||
|
|
|
|||
Loading…
Reference in New Issue