From 0d949e11ff8401f1a3e6f098dbf337d7b12788cc Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Mon, 5 Nov 2018 14:55:56 -0600 Subject: [PATCH] Support bitcasts of 16-bit types. --- .../comp/bitcast-16bit-1.invalid.comp | 22 ++++++++++ .../comp/bitcast-16bit-2.invalid.comp | 28 ++++++++++++ .../shaders/comp/bitcast-16bit-1.invalid.comp | 38 ++++++++++++++++ .../shaders/comp/bitcast-16bit-2.invalid.comp | 43 +++++++++++++++++++ .../comp/bitcast-16bit-1.invalid.comp | 24 +++++++++++ .../comp/bitcast-16bit-2.invalid.comp | 29 +++++++++++++ .../shaders/comp/bitcast-16bit-1.invalid.comp | 38 ++++++++++++++++ .../shaders/comp/bitcast-16bit-2.invalid.comp | 43 +++++++++++++++++++ shaders-msl/comp/bitcast-16bit-1.invalid.comp | 23 ++++++++++ shaders-msl/comp/bitcast-16bit-2.invalid.comp | 26 +++++++++++ shaders/comp/bitcast-16bit-1.invalid.comp | 23 ++++++++++ shaders/comp/bitcast-16bit-2.invalid.comp | 26 +++++++++++ spirv_glsl.cpp | 42 +++++++++++++++--- spirv_msl.cpp | 21 ++++++--- 14 files changed, 415 insertions(+), 11 deletions(-) create mode 100644 reference/opt/shaders-msl/comp/bitcast-16bit-1.invalid.comp create mode 100644 reference/opt/shaders-msl/comp/bitcast-16bit-2.invalid.comp create mode 100644 reference/opt/shaders/comp/bitcast-16bit-1.invalid.comp create mode 100644 reference/opt/shaders/comp/bitcast-16bit-2.invalid.comp create mode 100644 reference/shaders-msl/comp/bitcast-16bit-1.invalid.comp create mode 100644 reference/shaders-msl/comp/bitcast-16bit-2.invalid.comp create mode 100644 reference/shaders/comp/bitcast-16bit-1.invalid.comp create mode 100644 reference/shaders/comp/bitcast-16bit-2.invalid.comp create mode 100644 shaders-msl/comp/bitcast-16bit-1.invalid.comp create mode 100644 shaders-msl/comp/bitcast-16bit-2.invalid.comp create mode 100644 shaders/comp/bitcast-16bit-1.invalid.comp create mode 100644 shaders/comp/bitcast-16bit-2.invalid.comp diff --git a/reference/opt/shaders-msl/comp/bitcast-16bit-1.invalid.comp b/reference/opt/shaders-msl/comp/bitcast-16bit-1.invalid.comp new file mode 100644 index 00000000..5671cd7c --- /dev/null +++ b/reference/opt/shaders-msl/comp/bitcast-16bit-1.invalid.comp @@ -0,0 +1,22 @@ +#include +#include + +using namespace metal; + +struct SSBO0 +{ + short4 inputs[1]; +}; + +struct SSBO1 +{ + int4 outputs[1]; +}; + +kernel void main0(device SSBO0& _25 [[buffer(0)]], device SSBO1& _39 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + _39.outputs[gl_GlobalInvocationID.x].x = int(as_type(as_type(_25.inputs[gl_GlobalInvocationID.x].xy) + half2(half(1)))); + _39.outputs[gl_GlobalInvocationID.x].y = as_type(_25.inputs[gl_GlobalInvocationID.x].zw); + _39.outputs[gl_GlobalInvocationID.x].z = int(as_type(ushort2(_25.inputs[gl_GlobalInvocationID.x].xy))); +} + diff --git a/reference/opt/shaders-msl/comp/bitcast-16bit-2.invalid.comp b/reference/opt/shaders-msl/comp/bitcast-16bit-2.invalid.comp new file mode 100644 index 00000000..e7c20892 --- /dev/null +++ b/reference/opt/shaders-msl/comp/bitcast-16bit-2.invalid.comp @@ -0,0 +1,28 @@ +#include +#include + +using namespace metal; + +struct SSBO1 +{ + short4 outputs[1]; +}; + +struct SSBO0 +{ + int4 inputs[1]; +}; + +struct UBO +{ + half4 const0; +}; + +kernel void main0(device SSBO0& _29 [[buffer(0)]], device SSBO1& _21 [[buffer(1)]], constant UBO& _40 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + short2 _47 = as_type(_29.inputs[gl_GlobalInvocationID.x].x) + as_type(_40.const0.xy); + _21.outputs[gl_GlobalInvocationID.x] = short4(_47.x, _47.y, _21.outputs[gl_GlobalInvocationID.x].z, _21.outputs[gl_GlobalInvocationID.x].w); + short2 _66 = short2(as_type(uint(_29.inputs[gl_GlobalInvocationID.x].y)) - as_type(_40.const0.zw)); + _21.outputs[gl_GlobalInvocationID.x] = short4(_21.outputs[gl_GlobalInvocationID.x].x, _21.outputs[gl_GlobalInvocationID.x].y, _66.x, _66.y); +} + diff --git a/reference/opt/shaders/comp/bitcast-16bit-1.invalid.comp b/reference/opt/shaders/comp/bitcast-16bit-1.invalid.comp new file mode 100644 index 00000000..7420586f --- /dev/null +++ b/reference/opt/shaders/comp/bitcast-16bit-1.invalid.comp @@ -0,0 +1,38 @@ +#version 450 +#if defined(GL_AMD_gpu_shader_half_float) +#extension GL_AMD_gpu_shader_half_float : require +#elif defined(GL_NV_gpu_shader5) +#extension GL_NV_gpu_shader5 : require +#elif defined(GL_EXT_shader_16bit_storage) +#extension GL_EXT_shader_16bit_storage : require +#else +#error No extension available for FP16. +#endif +#if defined(GL_AMD_gpu_shader_int16) +#extension GL_AMD_gpu_shader_int16 : require +#elif defined(GL_EXT_shader_16bit_storage) +#extension GL_EXT_shader_16bit_storage : require +#else +#error No extension available for Int16. +#endif +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(binding = 0, std430) buffer SSBO0 +{ + i16vec4 inputs[]; +} _25; + +layout(binding = 1, std430) buffer SSBO1 +{ + ivec4 outputs[]; +} _39; + +void main() +{ + uint ident = gl_GlobalInvocationID.x; + f16vec2 a = int16BitsToFloat16(_25.inputs[ident].xy); + _39.outputs[ident].x = int(packFloat2x16(a + f16vec2(float16_t(1), float16_t(1)))); + _39.outputs[ident].y = packInt2x16(_25.inputs[ident].zw); + _39.outputs[ident].z = int(packUint2x16(u16vec2(_25.inputs[ident].xy))); +} + diff --git a/reference/opt/shaders/comp/bitcast-16bit-2.invalid.comp b/reference/opt/shaders/comp/bitcast-16bit-2.invalid.comp new file mode 100644 index 00000000..416bc4a2 --- /dev/null +++ b/reference/opt/shaders/comp/bitcast-16bit-2.invalid.comp @@ -0,0 +1,43 @@ +#version 450 +#if defined(GL_AMD_gpu_shader_int16) +#extension GL_AMD_gpu_shader_int16 : require +#elif defined(GL_EXT_shader_16bit_storage) +#extension GL_EXT_shader_16bit_storage : require +#else +#error No extension available for Int16. +#endif +#if defined(GL_AMD_gpu_shader_half_float) +#extension GL_AMD_gpu_shader_half_float : require +#elif defined(GL_NV_gpu_shader5) +#extension GL_NV_gpu_shader5 : require +#elif defined(GL_EXT_shader_16bit_storage) +#extension GL_EXT_shader_16bit_storage : require +#else +#error No extension available for FP16. +#endif +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(binding = 1, std430) buffer SSBO1 +{ + i16vec4 outputs[]; +} _21; + +layout(binding = 0, std430) buffer SSBO0 +{ + ivec4 inputs[]; +} _29; + +layout(binding = 2, std140) uniform UBO +{ + f16vec4 const0; +} _40; + +void main() +{ + uint ident = gl_GlobalInvocationID.x; + i16vec2 _47 = unpackInt2x16(_29.inputs[ident].x) + float16BitsToInt16(_40.const0.xy); + _21.outputs[ident] = i16vec4(_47.x, _47.y, _21.outputs[ident].z, _21.outputs[ident].w); + i16vec2 _66 = i16vec2(unpackUint2x16(uint(_29.inputs[ident].y)) - float16BitsToUint16(_40.const0.zw)); + _21.outputs[ident] = i16vec4(_21.outputs[ident].x, _21.outputs[ident].y, _66.x, _66.y); +} + diff --git a/reference/shaders-msl/comp/bitcast-16bit-1.invalid.comp b/reference/shaders-msl/comp/bitcast-16bit-1.invalid.comp new file mode 100644 index 00000000..ddf38a60 --- /dev/null +++ b/reference/shaders-msl/comp/bitcast-16bit-1.invalid.comp @@ -0,0 +1,24 @@ +#include +#include + +using namespace metal; + +struct SSBO0 +{ + short4 inputs[1]; +}; + +struct SSBO1 +{ + int4 outputs[1]; +}; + +kernel void main0(device SSBO0& _25 [[buffer(0)]], device SSBO1& _39 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + uint ident = gl_GlobalInvocationID.x; + half2 a = as_type(_25.inputs[ident].xy); + _39.outputs[ident].x = int(as_type(a + half2(half(1), half(1)))); + _39.outputs[ident].y = as_type(_25.inputs[ident].zw); + _39.outputs[ident].z = int(as_type(ushort2(_25.inputs[ident].xy))); +} + diff --git a/reference/shaders-msl/comp/bitcast-16bit-2.invalid.comp b/reference/shaders-msl/comp/bitcast-16bit-2.invalid.comp new file mode 100644 index 00000000..1dac3969 --- /dev/null +++ b/reference/shaders-msl/comp/bitcast-16bit-2.invalid.comp @@ -0,0 +1,29 @@ +#include +#include + +using namespace metal; + +struct SSBO1 +{ + short4 outputs[1]; +}; + +struct SSBO0 +{ + int4 inputs[1]; +}; + +struct UBO +{ + half4 const0; +}; + +kernel void main0(device SSBO0& _29 [[buffer(0)]], device SSBO1& _21 [[buffer(1)]], constant UBO& _40 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + uint ident = gl_GlobalInvocationID.x; + short2 _47 = as_type(_29.inputs[ident].x) + as_type(_40.const0.xy); + _21.outputs[ident] = short4(_47.x, _47.y, _21.outputs[ident].z, _21.outputs[ident].w); + short2 _66 = short2(as_type(uint(_29.inputs[ident].y)) - as_type(_40.const0.zw)); + _21.outputs[ident] = short4(_21.outputs[ident].x, _21.outputs[ident].y, _66.x, _66.y); +} + diff --git a/reference/shaders/comp/bitcast-16bit-1.invalid.comp b/reference/shaders/comp/bitcast-16bit-1.invalid.comp new file mode 100644 index 00000000..7420586f --- /dev/null +++ b/reference/shaders/comp/bitcast-16bit-1.invalid.comp @@ -0,0 +1,38 @@ +#version 450 +#if defined(GL_AMD_gpu_shader_half_float) +#extension GL_AMD_gpu_shader_half_float : require +#elif defined(GL_NV_gpu_shader5) +#extension GL_NV_gpu_shader5 : require +#elif defined(GL_EXT_shader_16bit_storage) +#extension GL_EXT_shader_16bit_storage : require +#else +#error No extension available for FP16. +#endif +#if defined(GL_AMD_gpu_shader_int16) +#extension GL_AMD_gpu_shader_int16 : require +#elif defined(GL_EXT_shader_16bit_storage) +#extension GL_EXT_shader_16bit_storage : require +#else +#error No extension available for Int16. +#endif +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(binding = 0, std430) buffer SSBO0 +{ + i16vec4 inputs[]; +} _25; + +layout(binding = 1, std430) buffer SSBO1 +{ + ivec4 outputs[]; +} _39; + +void main() +{ + uint ident = gl_GlobalInvocationID.x; + f16vec2 a = int16BitsToFloat16(_25.inputs[ident].xy); + _39.outputs[ident].x = int(packFloat2x16(a + f16vec2(float16_t(1), float16_t(1)))); + _39.outputs[ident].y = packInt2x16(_25.inputs[ident].zw); + _39.outputs[ident].z = int(packUint2x16(u16vec2(_25.inputs[ident].xy))); +} + diff --git a/reference/shaders/comp/bitcast-16bit-2.invalid.comp b/reference/shaders/comp/bitcast-16bit-2.invalid.comp new file mode 100644 index 00000000..416bc4a2 --- /dev/null +++ b/reference/shaders/comp/bitcast-16bit-2.invalid.comp @@ -0,0 +1,43 @@ +#version 450 +#if defined(GL_AMD_gpu_shader_int16) +#extension GL_AMD_gpu_shader_int16 : require +#elif defined(GL_EXT_shader_16bit_storage) +#extension GL_EXT_shader_16bit_storage : require +#else +#error No extension available for Int16. +#endif +#if defined(GL_AMD_gpu_shader_half_float) +#extension GL_AMD_gpu_shader_half_float : require +#elif defined(GL_NV_gpu_shader5) +#extension GL_NV_gpu_shader5 : require +#elif defined(GL_EXT_shader_16bit_storage) +#extension GL_EXT_shader_16bit_storage : require +#else +#error No extension available for FP16. +#endif +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(binding = 1, std430) buffer SSBO1 +{ + i16vec4 outputs[]; +} _21; + +layout(binding = 0, std430) buffer SSBO0 +{ + ivec4 inputs[]; +} _29; + +layout(binding = 2, std140) uniform UBO +{ + f16vec4 const0; +} _40; + +void main() +{ + uint ident = gl_GlobalInvocationID.x; + i16vec2 _47 = unpackInt2x16(_29.inputs[ident].x) + float16BitsToInt16(_40.const0.xy); + _21.outputs[ident] = i16vec4(_47.x, _47.y, _21.outputs[ident].z, _21.outputs[ident].w); + i16vec2 _66 = i16vec2(unpackUint2x16(uint(_29.inputs[ident].y)) - float16BitsToUint16(_40.const0.zw)); + _21.outputs[ident] = i16vec4(_21.outputs[ident].x, _21.outputs[ident].y, _66.x, _66.y); +} + diff --git a/shaders-msl/comp/bitcast-16bit-1.invalid.comp b/shaders-msl/comp/bitcast-16bit-1.invalid.comp new file mode 100644 index 00000000..0c21cda3 --- /dev/null +++ b/shaders-msl/comp/bitcast-16bit-1.invalid.comp @@ -0,0 +1,23 @@ +#version 450 core +#extension GL_AMD_gpu_shader_half_float : require +#extension GL_AMD_gpu_shader_int16 : require +layout(local_size_x = 1) in; + +layout(binding = 0, std430) buffer SSBO0 +{ + i16vec4 inputs[]; +}; + +layout(binding = 1, std430) buffer SSBO1 +{ + ivec4 outputs[]; +}; + +void main() +{ + uint ident = gl_GlobalInvocationID.x; + f16vec2 a = int16BitsToFloat16(inputs[ident].xy); + outputs[ident].x = int(packFloat2x16(a + f16vec2(1, 1))); + outputs[ident].y = packInt2x16(inputs[ident].zw); + outputs[ident].z = int(packUint2x16(u16vec2(inputs[ident].xy))); +} diff --git a/shaders-msl/comp/bitcast-16bit-2.invalid.comp b/shaders-msl/comp/bitcast-16bit-2.invalid.comp new file mode 100644 index 00000000..6bb66241 --- /dev/null +++ b/shaders-msl/comp/bitcast-16bit-2.invalid.comp @@ -0,0 +1,26 @@ +#version 450 core +#extension GL_AMD_gpu_shader_half_float : require +#extension GL_AMD_gpu_shader_int16 : require +layout(local_size_x = 1) in; + +layout(binding = 0, std430) buffer SSBO0 +{ + ivec4 inputs[]; +}; + +layout(binding = 1, std430) buffer SSBO1 +{ + i16vec4 outputs[]; +}; + +layout(binding = 2) uniform UBO +{ + f16vec4 const0; +}; + +void main() +{ + uint ident = gl_GlobalInvocationID.x; + outputs[ident].xy = unpackInt2x16(inputs[ident].x) + float16BitsToInt16(const0.xy); + outputs[ident].zw = i16vec2(unpackUint2x16(uint(inputs[ident].y)) - float16BitsToUint16(const0.zw)); +} diff --git a/shaders/comp/bitcast-16bit-1.invalid.comp b/shaders/comp/bitcast-16bit-1.invalid.comp new file mode 100644 index 00000000..0c21cda3 --- /dev/null +++ b/shaders/comp/bitcast-16bit-1.invalid.comp @@ -0,0 +1,23 @@ +#version 450 core +#extension GL_AMD_gpu_shader_half_float : require +#extension GL_AMD_gpu_shader_int16 : require +layout(local_size_x = 1) in; + +layout(binding = 0, std430) buffer SSBO0 +{ + i16vec4 inputs[]; +}; + +layout(binding = 1, std430) buffer SSBO1 +{ + ivec4 outputs[]; +}; + +void main() +{ + uint ident = gl_GlobalInvocationID.x; + f16vec2 a = int16BitsToFloat16(inputs[ident].xy); + outputs[ident].x = int(packFloat2x16(a + f16vec2(1, 1))); + outputs[ident].y = packInt2x16(inputs[ident].zw); + outputs[ident].z = int(packUint2x16(u16vec2(inputs[ident].xy))); +} diff --git a/shaders/comp/bitcast-16bit-2.invalid.comp b/shaders/comp/bitcast-16bit-2.invalid.comp new file mode 100644 index 00000000..6bb66241 --- /dev/null +++ b/shaders/comp/bitcast-16bit-2.invalid.comp @@ -0,0 +1,26 @@ +#version 450 core +#extension GL_AMD_gpu_shader_half_float : require +#extension GL_AMD_gpu_shader_int16 : require +layout(local_size_x = 1) in; + +layout(binding = 0, std430) buffer SSBO0 +{ + ivec4 inputs[]; +}; + +layout(binding = 1, std430) buffer SSBO1 +{ + i16vec4 outputs[]; +}; + +layout(binding = 2) uniform UBO +{ + f16vec4 const0; +}; + +void main() +{ + uint ident = gl_GlobalInvocationID.x; + outputs[ident].xy = unpackInt2x16(inputs[ident].x) + float16BitsToInt16(const0.xy); + outputs[ident].zw = i16vec2(unpackUint2x16(uint(inputs[ident].y)) - float16BitsToUint16(const0.zw)); +} diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index f2dd2826..3b88e5a5 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -1851,20 +1851,22 @@ void CompilerGLSL::replace_illegal_names() "dFdx", "dFdxCoarse", "dFdxFine", "dFdy", "dFdyCoarse", "dFdyFine", "distance", "dot", "EmitStreamVertex", "EmitVertex", "EndPrimitive", "EndStreamPrimitive", "equal", "exp", "exp2", - "faceforward", "findLSB", "findMSB", "floatBitsToInt", "floatBitsToUint", "floor", "fma", "fract", "frexp", "fwidth", "fwidthCoarse", "fwidthFine", + "faceforward", "findLSB", "findMSB", "float16BitsToInt16", "float16BitsToUint16", "floatBitsToInt", "floatBitsToUint", "floor", "fma", "fract", + "frexp", "fwidth", "fwidthCoarse", "fwidthFine", "greaterThan", "greaterThanEqual", "groupMemoryBarrier", "imageAtomicAdd", "imageAtomicAnd", "imageAtomicCompSwap", "imageAtomicExchange", "imageAtomicMax", "imageAtomicMin", "imageAtomicOr", "imageAtomicXor", - "imageLoad", "imageSamples", "imageSize", "imageStore", "imulExtended", "intBitsToFloat", "interpolateAtOffset", "interpolateAtCentroid", "interpolateAtSample", + "imageLoad", "imageSamples", "imageSize", "imageStore", "imulExtended", "int16BitsToFloat16", "intBitsToFloat", "interpolateAtOffset", "interpolateAtCentroid", "interpolateAtSample", "inverse", "inversesqrt", "isinf", "isnan", "ldexp", "length", "lessThan", "lessThanEqual", "log", "log2", "matrixCompMult", "max", "memoryBarrier", "memoryBarrierAtomicCounter", "memoryBarrierBuffer", "memoryBarrierImage", "memoryBarrierShared", "min", "mix", "mod", "modf", "noise", "noise1", "noise2", "noise3", "noise4", "normalize", "not", "notEqual", - "outerProduct", "packDouble2x32", "packHalf2x16", "packSnorm2x16", "packSnorm4x8", "packUnorm2x16", "packUnorm4x8", "pow", + "outerProduct", "packDouble2x32", "packHalf2x16", "packInt2x16", "packInt4x16", "packSnorm2x16", "packSnorm4x8", + "packUint2x16", "packUint4x16", "packUnorm2x16", "packUnorm4x8", "pow", "radians", "reflect", "refract", "round", "roundEven", "sign", "sin", "sinh", "smoothstep", "sqrt", "step", "tan", "tanh", "texelFetch", "texelFetchOffset", "texture", "textureGather", "textureGatherOffset", "textureGatherOffsets", "textureGrad", "textureGradOffset", "textureLod", "textureLodOffset", "textureOffset", "textureProj", "textureProjGrad", "textureProjGradOffset", "textureProjLod", "textureProjLodOffset", "textureProjOffset", "textureQueryLevels", "textureQueryLod", "textureSamples", "textureSize", - "transpose", "trunc", "uaddCarry", "uintBitsToFloat", "umulExtended", "unpackDouble2x32", "unpackHalf2x16", "unpackSnorm2x16", "unpackSnorm4x8", - "unpackUnorm2x16", "unpackUnorm4x8", "usubBorrow", + "transpose", "trunc", "uaddCarry", "uint16BitsToFloat16", "uintBitsToFloat", "umulExtended", "unpackDouble2x32", "unpackHalf2x16", "unpackInt2x16", "unpackInt4x16", + "unpackSnorm2x16", "unpackSnorm4x8", "unpackUint2x16", "unpackUint4x16", "unpackUnorm2x16", "unpackUnorm4x8", "usubBorrow", "active", "asm", "atomic_uint", "attribute", "bool", "break", "buffer", "bvec2", "bvec3", "bvec4", "case", "cast", "centroid", "class", "coherent", "common", "const", "continue", "default", "discard", @@ -5150,12 +5152,16 @@ case OpGroupNonUniform##op: \ string CompilerGLSL::bitcast_glsl_op(const SPIRType &out_type, const SPIRType &in_type) { - if (out_type.basetype == SPIRType::UInt && in_type.basetype == SPIRType::Int) + if (out_type.basetype == SPIRType::UShort && in_type.basetype == SPIRType::Short) + return type_to_glsl(out_type); + else if (out_type.basetype == SPIRType::UInt && in_type.basetype == SPIRType::Int) return type_to_glsl(out_type); else if (out_type.basetype == SPIRType::UInt64 && in_type.basetype == SPIRType::Int64) return type_to_glsl(out_type); else if (out_type.basetype == SPIRType::UInt && in_type.basetype == SPIRType::Float) return "floatBitsToUint"; + else if (out_type.basetype == SPIRType::Short && in_type.basetype == SPIRType::UShort) + return type_to_glsl(out_type); else if (out_type.basetype == SPIRType::Int && in_type.basetype == SPIRType::UInt) return type_to_glsl(out_type); else if (out_type.basetype == SPIRType::Int64 && in_type.basetype == SPIRType::UInt64) @@ -5174,12 +5180,36 @@ string CompilerGLSL::bitcast_glsl_op(const SPIRType &out_type, const SPIRType &i return "int64BitsToDouble"; else if (out_type.basetype == SPIRType::Double && in_type.basetype == SPIRType::UInt64) return "uint64BitsToDouble"; + else if (out_type.basetype == SPIRType::Short && in_type.basetype == SPIRType::Half) + return "float16BitsToInt16"; + else if (out_type.basetype == SPIRType::UShort && in_type.basetype == SPIRType::Half) + return "float16BitsToUint16"; + else if (out_type.basetype == SPIRType::Half && in_type.basetype == SPIRType::Short) + return "int16BitsToFloat16"; + else if (out_type.basetype == SPIRType::Half && in_type.basetype == SPIRType::UShort) + return "uint16BitsToFloat16"; else if (out_type.basetype == SPIRType::UInt64 && in_type.basetype == SPIRType::UInt && in_type.vecsize == 2) return "packUint2x32"; else if (out_type.basetype == SPIRType::Half && in_type.basetype == SPIRType::UInt && in_type.vecsize == 1) return "unpackFloat2x16"; else if (out_type.basetype == SPIRType::UInt && in_type.basetype == SPIRType::Half && in_type.vecsize == 2) return "packFloat2x16"; + else if (out_type.basetype == SPIRType::Int && in_type.basetype == SPIRType::Short && in_type.vecsize == 2) + return "packInt2x16"; + else if (out_type.basetype == SPIRType::Short && in_type.basetype == SPIRType::Int && in_type.vecsize == 1) + return "unpackInt2x16"; + else if (out_type.basetype == SPIRType::UInt && in_type.basetype == SPIRType::UShort && in_type.vecsize == 2) + return "packUint2x16"; + else if (out_type.basetype == SPIRType::UShort && in_type.basetype == SPIRType::UInt && in_type.vecsize == 1) + return "unpackUint2x16"; + else if (out_type.basetype == SPIRType::Int64 && in_type.basetype == SPIRType::Short && in_type.vecsize == 4) + return "packInt4x16"; + else if (out_type.basetype == SPIRType::Short && in_type.basetype == SPIRType::Int64 && in_type.vecsize == 1) + return "unpackInt4x16"; + else if (out_type.basetype == SPIRType::UInt64 && in_type.basetype == SPIRType::UShort && in_type.vecsize == 4) + return "packUint4x16"; + else if (out_type.basetype == SPIRType::UShort && in_type.basetype == SPIRType::UInt64 && in_type.vecsize == 1) + return "unpackUint4x16"; else return ""; } diff --git a/spirv_msl.cpp b/spirv_msl.cpp index 001636c7..589914c9 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -983,9 +983,8 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage) mbr_idx++; } } - else if (type.basetype == SPIRType::Boolean || type.basetype == SPIRType::Char || - type_is_integral(type) || type_is_floating_point(type) || - type.basetype == SPIRType::Boolean) + else if (type.basetype == SPIRType::Boolean || type.basetype == SPIRType::Char || type_is_integral(type) || + type_is_floating_point(type) || type.basetype == SPIRType::Boolean) { bool is_builtin = is_builtin_variable(*p_var); BuiltIn builtin = BuiltIn(get_decoration(p_var->self, DecorationBuiltIn)); @@ -4501,7 +4500,9 @@ string CompilerMSL::image_type_glsl(const SPIRType &type, uint32_t id) string CompilerMSL::bitcast_glsl_op(const SPIRType &out_type, const SPIRType &in_type) { - if ((out_type.basetype == SPIRType::UInt && in_type.basetype == SPIRType::Int) || + if ((out_type.basetype == SPIRType::UShort && in_type.basetype == SPIRType::Short) || + (out_type.basetype == SPIRType::Short && in_type.basetype == SPIRType::UShort) || + (out_type.basetype == SPIRType::UInt && in_type.basetype == SPIRType::Int) || (out_type.basetype == SPIRType::Int && in_type.basetype == SPIRType::UInt) || (out_type.basetype == SPIRType::UInt64 && in_type.basetype == SPIRType::Int64) || (out_type.basetype == SPIRType::Int64 && in_type.basetype == SPIRType::UInt64)) @@ -4516,7 +4517,17 @@ string CompilerMSL::bitcast_glsl_op(const SPIRType &out_type, const SPIRType &in (out_type.basetype == SPIRType::Double && in_type.basetype == SPIRType::Int64) || (out_type.basetype == SPIRType::Double && in_type.basetype == SPIRType::UInt64) || (out_type.basetype == SPIRType::Half && in_type.basetype == SPIRType::UInt) || - (out_type.basetype == SPIRType::UInt && in_type.basetype == SPIRType::Half)) + (out_type.basetype == SPIRType::UInt && in_type.basetype == SPIRType::Half) || + (out_type.basetype == SPIRType::Half && in_type.basetype == SPIRType::Int) || + (out_type.basetype == SPIRType::Int && in_type.basetype == SPIRType::Half) || + (out_type.basetype == SPIRType::Half && in_type.basetype == SPIRType::UShort) || + (out_type.basetype == SPIRType::UShort && in_type.basetype == SPIRType::Half) || + (out_type.basetype == SPIRType::Half && in_type.basetype == SPIRType::Short) || + (out_type.basetype == SPIRType::Short && in_type.basetype == SPIRType::Half) || + (out_type.basetype == SPIRType::UInt && in_type.basetype == SPIRType::UShort && in_type.vecsize == 2) || + (out_type.basetype == SPIRType::UShort && in_type.basetype == SPIRType::UInt && in_type.vecsize == 1) || + (out_type.basetype == SPIRType::Int && in_type.basetype == SPIRType::Short && in_type.vecsize == 2) || + (out_type.basetype == SPIRType::Short && in_type.basetype == SPIRType::Int && in_type.vecsize == 1)) return "as_type<" + type_to_glsl(out_type) + ">"; return "";