Deal correctly with sign on bitfield operations.

Need a lot of special purpose implementation functions for these.

reference/shaders-hlsl-no-opt/asm/comp/bitfield-signed-operations.asm.comp

@@ -0,0 +1,105 @@
+RWByteAddressBuffer _3 : register(u0);
+
+uint SPIRV_Cross_bitfieldInsert(uint Base, uint Insert, uint Offset, uint Count)
+{
+    uint Mask = Count == 32 ? 0xffffffff : (((1u << Count) - 1) << (Offset & 31));
+    return (Base & ~Mask) | ((Insert << Offset) & Mask);
+}
+
+uint2 SPIRV_Cross_bitfieldInsert(uint2 Base, uint2 Insert, uint Offset, uint Count)
+{
+    uint Mask = Count == 32 ? 0xffffffff : (((1u << Count) - 1) << (Offset & 31));
+    return (Base & ~Mask) | ((Insert << Offset) & Mask);
+}
+
+uint3 SPIRV_Cross_bitfieldInsert(uint3 Base, uint3 Insert, uint Offset, uint Count)
+{
+    uint Mask = Count == 32 ? 0xffffffff : (((1u << Count) - 1) << (Offset & 31));
+    return (Base & ~Mask) | ((Insert << Offset) & Mask);
+}
+
+uint4 SPIRV_Cross_bitfieldInsert(uint4 Base, uint4 Insert, uint Offset, uint Count)
+{
+    uint Mask = Count == 32 ? 0xffffffff : (((1u << Count) - 1) << (Offset & 31));
+    return (Base & ~Mask) | ((Insert << Offset) & Mask);
+}
+
+uint SPIRV_Cross_bitfieldUExtract(uint Base, uint Offset, uint Count)
+{
+    uint Mask = Count == 32 ? 0xffffffff : ((1 << Count) - 1);
+    return (Base >> Offset) & Mask;
+}
+
+uint2 SPIRV_Cross_bitfieldUExtract(uint2 Base, uint Offset, uint Count)
+{
+    uint Mask = Count == 32 ? 0xffffffff : ((1 << Count) - 1);
+    return (Base >> Offset) & Mask;
+}
+
+uint3 SPIRV_Cross_bitfieldUExtract(uint3 Base, uint Offset, uint Count)
+{
+    uint Mask = Count == 32 ? 0xffffffff : ((1 << Count) - 1);
+    return (Base >> Offset) & Mask;
+}
+
+uint4 SPIRV_Cross_bitfieldUExtract(uint4 Base, uint Offset, uint Count)
+{
+    uint Mask = Count == 32 ? 0xffffffff : ((1 << Count) - 1);
+    return (Base >> Offset) & Mask;
+}
+
+int SPIRV_Cross_bitfieldSExtract(int Base, int Offset, int Count)
+{
+    int Mask = Count == 32 ? -1 : ((1 << Count) - 1);
+    int Masked = (Base >> Offset) & Mask;
+    int ExtendShift = (32 - Count) & 31;
+    return (Masked << ExtendShift) >> ExtendShift;
+}
+
+int2 SPIRV_Cross_bitfieldSExtract(int2 Base, int Offset, int Count)
+{
+    int Mask = Count == 32 ? -1 : ((1 << Count) - 1);
+    int2 Masked = (Base >> Offset) & Mask;
+    int ExtendShift = (32 - Count) & 31;
+    return (Masked << ExtendShift) >> ExtendShift;
+}
+
+int3 SPIRV_Cross_bitfieldSExtract(int3 Base, int Offset, int Count)
+{
+    int Mask = Count == 32 ? -1 : ((1 << Count) - 1);
+    int3 Masked = (Base >> Offset) & Mask;
+    int ExtendShift = (32 - Count) & 31;
+    return (Masked << ExtendShift) >> ExtendShift;
+}
+
+int4 SPIRV_Cross_bitfieldSExtract(int4 Base, int Offset, int Count)
+{
+    int Mask = Count == 32 ? -1 : ((1 << Count) - 1);
+    int4 Masked = (Base >> Offset) & Mask;
+    int ExtendShift = (32 - Count) & 31;
+    return (Masked << ExtendShift) >> ExtendShift;
+}
+
+void comp_main()
+{
+    int4 _19 = int4(_3.Load4(0));
+    uint4 _20 = _3.Load4(16);
+    _3.Store4(0, uint4(countbits(_19)));
+    _3.Store4(16, uint4(countbits(_19)));
+    _3.Store4(0, uint4(int4(countbits(_20))));
+    _3.Store4(16, countbits(_20));
+    _3.Store4(0, uint4(reversebits(_19)));
+    _3.Store4(16, reversebits(_20));
+    _3.Store4(0, uint4(SPIRV_Cross_bitfieldSExtract(_19, 1, 11u)));
+    _3.Store4(16, SPIRV_Cross_bitfieldSExtract(_20, 11u, 1));
+    _3.Store4(0, uint4(SPIRV_Cross_bitfieldUExtract(_19, 1, 11u)));
+    _3.Store4(16, SPIRV_Cross_bitfieldUExtract(_20, 11u, 1));
+    _3.Store4(0, uint4(int4(SPIRV_Cross_bitfieldInsert(_19, _19.wzyx, 1, 11u))));
+    _3.Store4(16, SPIRV_Cross_bitfieldInsert(_20, _20.wzyx, 11u, 1));
+}
+
+[numthreads(1, 1, 1)]
+void main()
+{
+    comp_main();
+}

reference/shaders-hlsl-no-opt/comp/bitfield.comp

@@ -90,7 +90,7 @@ void comp_main()
     u = SPIRV_Cross_bitfieldInsert(u, 60u, 5, 4);
     u = reversebits(u);
     s = reversebits(s);
-    int v0 = countbits(u);
+    int v0 = int(countbits(u));
     int v1 = countbits(s);
     int v2 = int(firstbithigh(u));
     int v3 = firstbitlow(s);
@@ -100,7 +100,7 @@ void comp_main()
     u_1 = SPIRV_Cross_bitfieldInsert(u_1, uint3(60u, 60u, 60u), 5, 4);
     u_1 = reversebits(u_1);
     s_1 = reversebits(s_1);
-    int3 v0_1 = countbits(u_1);
+    int3 v0_1 = int3(countbits(u_1));
     int3 v1_1 = countbits(s_1);
     int3 v2_1 = int3(firstbithigh(u_1));
     int3 v3_1 = firstbitlow(s_1);

reference/shaders-msl-no-opt/asm/comp/bitfield-signed-operations.asm.comp

@@ -0,0 +1,29 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct SSBO
+{
+    int4 ints;
+    uint4 uints;
+};
+
+kernel void main0(device SSBO& _3 [[buffer(0)]])
+{
+    int4 _19 = _3.ints;
+    uint4 _20 = _3.uints;
+    _3.ints = popcount(_19);
+    _3.uints = uint4(popcount(_19));
+    _3.ints = int4(popcount(_20));
+    _3.uints = popcount(_20);
+    _3.ints = reverse_bits(_19);
+    _3.uints = reverse_bits(_20);
+    _3.ints = extract_bits(_19, uint(1), 11u);
+    _3.uints = uint4(extract_bits(int4(_20), 11u, uint(1)));
+    _3.ints = int4(extract_bits(uint4(_19), uint(1), 11u));
+    _3.uints = extract_bits(_20, 11u, uint(1));
+    _3.ints = insert_bits(_19, _19.wzyx, uint(1), 11u);
+    _3.uints = insert_bits(_20, _20.wzyx, 11u, uint(1));
+}
+

reference/shaders-msl-no-opt/comp/bitfield.comp

@@ -31,13 +31,13 @@ kernel void main0()
 {
     int signed_value = 0;
     uint unsigned_value = 0u;
-    int s = extract_bits(signed_value, 5, 20);
-    uint u = extract_bits(unsigned_value, 6, 21);
-    s = insert_bits(s, 40, 5, 4);
-    u = insert_bits(u, 60u, 5, 4);
+    int s = extract_bits(signed_value, uint(5), uint(20));
+    uint u = extract_bits(unsigned_value, uint(6), uint(21));
+    s = insert_bits(s, 40, uint(5), uint(4));
+    u = insert_bits(u, 60u, uint(5), uint(4));
     u = reverse_bits(u);
     s = reverse_bits(s);
-    int v0 = popcount(u);
+    int v0 = int(popcount(u));
     int v1 = popcount(s);
     int v2 = int(spvFindUMSB(u));
     int v3 = spvFindSMSB(s);

reference/shaders-no-opt/asm/comp/bitfield-signed-operations.asm.comp

@@ -0,0 +1,27 @@
+#version 450
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(binding = 0, std430) buffer SSBO
+{
+    ivec4 ints;
+    uvec4 uints;
+} _3;
+
+void main()
+{
+    ivec4 _19 = _3.ints;
+    uvec4 _20 = _3.uints;
+    _3.ints = bitCount(_19);
+    _3.uints = uvec4(bitCount(_19));
+    _3.ints = bitCount(_20);
+    _3.uints = uvec4(bitCount(_20));
+    _3.ints = bitfieldReverse(_19);
+    _3.uints = bitfieldReverse(_20);
+    _3.ints = bitfieldExtract(_19, 1, int(11u));
+    _3.uints = uvec4(bitfieldExtract(ivec4(_20), int(11u), 1));
+    _3.ints = ivec4(bitfieldExtract(uvec4(_19), 1, int(11u)));
+    _3.uints = bitfieldExtract(_20, int(11u), 1);
+    _3.ints = bitfieldInsert(_19, _19.wzyx, 1, int(11u));
+    _3.uints = bitfieldInsert(_20, _20.wzyx, int(11u), 1);
+}
+

shaders-hlsl-no-opt/asm/comp/bitfield-signed-operations.asm.comp

@@ -0,0 +1,97 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 7
+; Bound: 26
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpSource GLSL 450
+               OpName %main "main"
+               OpName %SSBO "SSBO"
+               OpMemberName %SSBO 0 "ints"
+               OpMemberName %SSBO 1 "uints"
+               OpName %_ ""
+               OpMemberDecorate %SSBO 0 Offset 0
+               OpMemberDecorate %SSBO 1 Offset 16
+               OpDecorate %SSBO BufferBlock
+               OpDecorate %_ DescriptorSet 0
+               OpDecorate %_ Binding 0
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v4int = OpTypeVector %int 4
+       %uint = OpTypeInt 32 0
+     %v4uint = OpTypeVector %uint 4
+
+	 %int_1 = OpConstant %int 1
+	 %uint_11 = OpConstant %uint 11
+
+       %SSBO = OpTypeStruct %v4int %v4uint
+%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
+          %_ = OpVariable %_ptr_Uniform_SSBO Uniform
+      %int_0 = OpConstant %int 0
+%_ptr_Uniform_v4int = OpTypePointer Uniform %v4int
+%_ptr_Uniform_v4uint = OpTypePointer Uniform %v4uint
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+         %ints_ptr = OpAccessChain %_ptr_Uniform_v4int %_ %int_0
+         %uints_ptr = OpAccessChain %_ptr_Uniform_v4uint %_ %int_1
+         %ints = OpLoad %v4int %ints_ptr
+         %uints = OpLoad %v4uint %uints_ptr
+
+		 %ints_alt = OpVectorShuffle %v4int %ints %ints 3 2 1 0
+		 %uints_alt = OpVectorShuffle %v4uint %uints %uints 3 2 1 0
+
+         %int_to_int_popcount = OpBitCount %v4int %ints
+         %int_to_uint_popcount = OpBitCount %v4uint %ints
+         %uint_to_int_popcount = OpBitCount %v4int %uints
+         %uint_to_uint_popcount = OpBitCount %v4uint %uints
+
+		; BitReverse must have matching types w.r.t. sign, yay.
+         %int_to_int_reverse = OpBitReverse %v4int %ints
+         ;%int_to_uint_reverse = OpBitReverse %v4uint %ints
+         ;%uint_to_int_reverse = OpBitReverse %v4int %uints
+         %uint_to_uint_reverse = OpBitReverse %v4uint %uints
+
+		; Base and Result must match.
+         %int_to_int_sbit = OpBitFieldSExtract %v4int %ints %int_1 %uint_11
+         ;%int_to_uint_sbit = OpBitFieldSExtract %v4uint %ints %offset %count
+         ;%uint_to_int_sbit = OpBitFieldSExtract %v4int %uints %offset %count
+         %uint_to_uint_sbit = OpBitFieldSExtract %v4uint %uints %uint_11 %int_1
+
+		; Base and Result must match.
+         %int_to_int_ubit = OpBitFieldUExtract %v4int %ints %int_1 %uint_11
+         ;%int_to_uint_ubit = OpBitFieldUExtract %v4uint %ints %offset %count
+         ;%uint_to_int_ubit = OpBitFieldUExtract %v4int %uints %offset %count
+         %uint_to_uint_ubit = OpBitFieldUExtract %v4uint %uints %uint_11 %int_1
+
+		 %int_to_int_insert = OpBitFieldInsert %v4int %ints %ints_alt %int_1 %uint_11
+		 %uint_to_uint_insert = OpBitFieldInsert %v4uint %uints %uints_alt %uint_11 %int_1
+
+               OpStore %ints_ptr %int_to_int_popcount
+               OpStore %uints_ptr %int_to_uint_popcount
+               OpStore %ints_ptr %uint_to_int_popcount
+               OpStore %uints_ptr %uint_to_uint_popcount
+
+               OpStore %ints_ptr %int_to_int_reverse
+               ;OpStore %uints_ptr %int_to_uint_reverse
+               ;OpStore %ints_ptr %uint_to_int_reverse
+               OpStore %uints_ptr %uint_to_uint_reverse
+
+               OpStore %ints_ptr %int_to_int_sbit
+               ;OpStore %uints_ptr %int_to_uint_sbit
+               ;OpStore %ints_ptr %uint_to_int_sbit
+               OpStore %uints_ptr %uint_to_uint_sbit
+
+               OpStore %ints_ptr %int_to_int_ubit
+               ;OpStore %uints_ptr %int_to_uint_ubit
+               ;OpStore %ints_ptr %uint_to_int_ubit
+               OpStore %uints_ptr %uint_to_uint_ubit
+
+			   OpStore %ints_ptr %int_to_int_insert
+			   OpStore %uints_ptr %uint_to_uint_insert
+
+               OpReturn
+               OpFunctionEnd

shaders-msl-no-opt/asm/comp/bitfield-signed-operations.asm.comp

@@ -0,0 +1,97 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 7
+; Bound: 26
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpSource GLSL 450
+               OpName %main "main"
+               OpName %SSBO "SSBO"
+               OpMemberName %SSBO 0 "ints"
+               OpMemberName %SSBO 1 "uints"
+               OpName %_ ""
+               OpMemberDecorate %SSBO 0 Offset 0
+               OpMemberDecorate %SSBO 1 Offset 16
+               OpDecorate %SSBO BufferBlock
+               OpDecorate %_ DescriptorSet 0
+               OpDecorate %_ Binding 0
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v4int = OpTypeVector %int 4
+       %uint = OpTypeInt 32 0
+     %v4uint = OpTypeVector %uint 4
+
+	 %int_1 = OpConstant %int 1
+	 %uint_11 = OpConstant %uint 11
+
+       %SSBO = OpTypeStruct %v4int %v4uint
+%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
+          %_ = OpVariable %_ptr_Uniform_SSBO Uniform
+      %int_0 = OpConstant %int 0
+%_ptr_Uniform_v4int = OpTypePointer Uniform %v4int
+%_ptr_Uniform_v4uint = OpTypePointer Uniform %v4uint
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+         %ints_ptr = OpAccessChain %_ptr_Uniform_v4int %_ %int_0
+         %uints_ptr = OpAccessChain %_ptr_Uniform_v4uint %_ %int_1
+         %ints = OpLoad %v4int %ints_ptr
+         %uints = OpLoad %v4uint %uints_ptr
+
+		 %ints_alt = OpVectorShuffle %v4int %ints %ints 3 2 1 0
+		 %uints_alt = OpVectorShuffle %v4uint %uints %uints 3 2 1 0
+
+         %int_to_int_popcount = OpBitCount %v4int %ints
+         %int_to_uint_popcount = OpBitCount %v4uint %ints
+         %uint_to_int_popcount = OpBitCount %v4int %uints
+         %uint_to_uint_popcount = OpBitCount %v4uint %uints
+
+		; BitReverse must have matching types w.r.t. sign, yay.
+         %int_to_int_reverse = OpBitReverse %v4int %ints
+         ;%int_to_uint_reverse = OpBitReverse %v4uint %ints
+         ;%uint_to_int_reverse = OpBitReverse %v4int %uints
+         %uint_to_uint_reverse = OpBitReverse %v4uint %uints
+
+		; Base and Result must match.
+         %int_to_int_sbit = OpBitFieldSExtract %v4int %ints %int_1 %uint_11
+         ;%int_to_uint_sbit = OpBitFieldSExtract %v4uint %ints %offset %count
+         ;%uint_to_int_sbit = OpBitFieldSExtract %v4int %uints %offset %count
+         %uint_to_uint_sbit = OpBitFieldSExtract %v4uint %uints %uint_11 %int_1
+
+		; Base and Result must match.
+         %int_to_int_ubit = OpBitFieldUExtract %v4int %ints %int_1 %uint_11
+         ;%int_to_uint_ubit = OpBitFieldUExtract %v4uint %ints %offset %count
+         ;%uint_to_int_ubit = OpBitFieldUExtract %v4int %uints %offset %count
+         %uint_to_uint_ubit = OpBitFieldUExtract %v4uint %uints %uint_11 %int_1
+
+		 %int_to_int_insert = OpBitFieldInsert %v4int %ints %ints_alt %int_1 %uint_11
+		 %uint_to_uint_insert = OpBitFieldInsert %v4uint %uints %uints_alt %uint_11 %int_1
+
+               OpStore %ints_ptr %int_to_int_popcount
+               OpStore %uints_ptr %int_to_uint_popcount
+               OpStore %ints_ptr %uint_to_int_popcount
+               OpStore %uints_ptr %uint_to_uint_popcount
+
+               OpStore %ints_ptr %int_to_int_reverse
+               ;OpStore %uints_ptr %int_to_uint_reverse
+               ;OpStore %ints_ptr %uint_to_int_reverse
+               OpStore %uints_ptr %uint_to_uint_reverse
+
+               OpStore %ints_ptr %int_to_int_sbit
+               ;OpStore %uints_ptr %int_to_uint_sbit
+               ;OpStore %ints_ptr %uint_to_int_sbit
+               OpStore %uints_ptr %uint_to_uint_sbit
+
+               OpStore %ints_ptr %int_to_int_ubit
+               ;OpStore %uints_ptr %int_to_uint_ubit
+               ;OpStore %ints_ptr %uint_to_int_ubit
+               OpStore %uints_ptr %uint_to_uint_ubit
+
+			   OpStore %ints_ptr %int_to_int_insert
+			   OpStore %uints_ptr %uint_to_uint_insert
+
+               OpReturn
+               OpFunctionEnd

shaders-no-opt/asm/comp/bitfield-signed-operations.asm.comp

@@ -0,0 +1,97 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 7
+; Bound: 26
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpSource GLSL 450
+               OpName %main "main"
+               OpName %SSBO "SSBO"
+               OpMemberName %SSBO 0 "ints"
+               OpMemberName %SSBO 1 "uints"
+               OpName %_ ""
+               OpMemberDecorate %SSBO 0 Offset 0
+               OpMemberDecorate %SSBO 1 Offset 16
+               OpDecorate %SSBO BufferBlock
+               OpDecorate %_ DescriptorSet 0
+               OpDecorate %_ Binding 0
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v4int = OpTypeVector %int 4
+       %uint = OpTypeInt 32 0
+     %v4uint = OpTypeVector %uint 4
+
+	 %int_1 = OpConstant %int 1
+	 %uint_11 = OpConstant %uint 11
+
+       %SSBO = OpTypeStruct %v4int %v4uint
+%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
+          %_ = OpVariable %_ptr_Uniform_SSBO Uniform
+      %int_0 = OpConstant %int 0
+%_ptr_Uniform_v4int = OpTypePointer Uniform %v4int
+%_ptr_Uniform_v4uint = OpTypePointer Uniform %v4uint
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+         %ints_ptr = OpAccessChain %_ptr_Uniform_v4int %_ %int_0
+         %uints_ptr = OpAccessChain %_ptr_Uniform_v4uint %_ %int_1
+         %ints = OpLoad %v4int %ints_ptr
+         %uints = OpLoad %v4uint %uints_ptr
+
+		 %ints_alt = OpVectorShuffle %v4int %ints %ints 3 2 1 0
+		 %uints_alt = OpVectorShuffle %v4uint %uints %uints 3 2 1 0
+
+         %int_to_int_popcount = OpBitCount %v4int %ints
+         %int_to_uint_popcount = OpBitCount %v4uint %ints
+         %uint_to_int_popcount = OpBitCount %v4int %uints
+         %uint_to_uint_popcount = OpBitCount %v4uint %uints
+
+		; BitReverse must have matching types w.r.t. sign, yay.
+         %int_to_int_reverse = OpBitReverse %v4int %ints
+         ;%int_to_uint_reverse = OpBitReverse %v4uint %ints
+         ;%uint_to_int_reverse = OpBitReverse %v4int %uints
+         %uint_to_uint_reverse = OpBitReverse %v4uint %uints
+
+		; Base and Result must match.
+         %int_to_int_sbit = OpBitFieldSExtract %v4int %ints %int_1 %uint_11
+         ;%int_to_uint_sbit = OpBitFieldSExtract %v4uint %ints %offset %count
+         ;%uint_to_int_sbit = OpBitFieldSExtract %v4int %uints %offset %count
+         %uint_to_uint_sbit = OpBitFieldSExtract %v4uint %uints %uint_11 %int_1
+
+		; Base and Result must match.
+         %int_to_int_ubit = OpBitFieldUExtract %v4int %ints %int_1 %uint_11
+         ;%int_to_uint_ubit = OpBitFieldUExtract %v4uint %ints %offset %count
+         ;%uint_to_int_ubit = OpBitFieldUExtract %v4int %uints %offset %count
+         %uint_to_uint_ubit = OpBitFieldUExtract %v4uint %uints %uint_11 %int_1
+
+		 %int_to_int_insert = OpBitFieldInsert %v4int %ints %ints_alt %int_1 %uint_11
+		 %uint_to_uint_insert = OpBitFieldInsert %v4uint %uints %uints_alt %uint_11 %int_1
+
+               OpStore %ints_ptr %int_to_int_popcount
+               OpStore %uints_ptr %int_to_uint_popcount
+               OpStore %ints_ptr %uint_to_int_popcount
+               OpStore %uints_ptr %uint_to_uint_popcount
+
+               OpStore %ints_ptr %int_to_int_reverse
+               ;OpStore %uints_ptr %int_to_uint_reverse
+               ;OpStore %ints_ptr %uint_to_int_reverse
+               OpStore %uints_ptr %uint_to_uint_reverse
+
+               OpStore %ints_ptr %int_to_int_sbit
+               ;OpStore %uints_ptr %int_to_uint_sbit
+               ;OpStore %ints_ptr %uint_to_int_sbit
+               OpStore %uints_ptr %uint_to_uint_sbit
+
+               OpStore %ints_ptr %int_to_int_ubit
+               ;OpStore %uints_ptr %int_to_uint_ubit
+               ;OpStore %ints_ptr %uint_to_int_ubit
+               OpStore %uints_ptr %uint_to_uint_ubit
+
+			   OpStore %ints_ptr %int_to_int_insert
+			   OpStore %uints_ptr %uint_to_uint_insert
+
+               OpReturn
+               OpFunctionEnd

spirv_glsl.cpp

@@ -4225,6 +4225,57 @@ void CompilerGLSL::emit_unary_func_op_cast(uint32_t result_type, uint32_t result
 	inherit_expression_dependencies(result_id, op0);
 }
 
+// Very special case. Handling bitfieldExtract requires us to deal with different bitcasts of different signs
+// and different vector sizes all at once. Need a special purpose method here.
+void CompilerGLSL::emit_trinary_func_op_bitextract(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1,
+                                                   uint32_t op2, const char *op,
+                                                   SPIRType::BaseType expected_result_type,
+                                                   SPIRType::BaseType input_type0,
+                                                   SPIRType::BaseType input_type1,
+                                                   SPIRType::BaseType input_type2)
+{
+	auto &out_type = get<SPIRType>(result_type);
+	auto expected_type = out_type;
+	expected_type.basetype = input_type0;
+
+	string cast_op0 =
+			expression_type(op0).basetype != input_type0 ? bitcast_glsl(expected_type, op0) : to_unpacked_expression(op0);
+
+	auto op1_expr = to_unpacked_expression(op1);
+	auto op2_expr = to_unpacked_expression(op2);
+
+	// Use value casts here instead. Input must be exactly int or uint, but SPIR-V might be 16-bit.
+	expected_type.basetype = input_type1;
+	expected_type.vecsize = 1;
+	string cast_op1 =
+			expression_type(op1).basetype != input_type1 ? join(type_to_glsl_constructor(expected_type), "(", op1_expr, ")") : op1_expr;
+
+	expected_type.basetype = input_type2;
+	expected_type.vecsize = 1;
+	string cast_op2 =
+			expression_type(op2).basetype != input_type2 ? join(type_to_glsl_constructor(expected_type), "(", op2_expr, ")") : op2_expr;
+
+	string expr;
+	if (out_type.basetype != expected_result_type)
+	{
+		expected_type.vecsize = out_type.vecsize;
+		expected_type.basetype = expected_result_type;
+		expr = bitcast_glsl_op(out_type, expected_type);
+		expr += '(';
+		expr += join(op, "(", cast_op0, ", ", cast_op1, ", ", cast_op2, ")");
+		expr += ')';
+	}
+	else
+	{
+		expr += join(op, "(", cast_op0, ", ", cast_op1, ", ", cast_op2, ")");
+	}
+
+	emit_op(result_type, result_id, expr, should_forward(op0) && should_forward(op1) && should_forward(op2));
+	inherit_expression_dependencies(result_id, op0);
+	inherit_expression_dependencies(result_id, op1);
+	inherit_expression_dependencies(result_id, op2);
+}
+
 void CompilerGLSL::emit_trinary_func_op_cast(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1,
                                              uint32_t op2, const char *op, SPIRType::BaseType input_type)
 {
@@ -4313,6 +4364,46 @@ void CompilerGLSL::emit_quaternary_func_op(uint32_t result_type, uint32_t result
 	inherit_expression_dependencies(result_id, op3);
 }
 
+void CompilerGLSL::emit_bitfield_insert_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1,
+                                           uint32_t op2, uint32_t op3, const char *op,
+                                           SPIRType::BaseType offset_count_type)
+{
+	// Only need to cast offset/count arguments. Types of base/insert must be same as result type,
+	// and bitfieldInsert is sign invariant.
+	bool forward = should_forward(op0) && should_forward(op1) &&
+	               should_forward(op2) && should_forward(op3);
+
+	auto op0_expr = to_unpacked_expression(op0);
+	auto op1_expr = to_unpacked_expression(op1);
+	auto op2_expr = to_unpacked_expression(op2);
+	auto op3_expr = to_unpacked_expression(op3);
+
+	SPIRType target_type;
+	target_type.vecsize = 1;
+	target_type.basetype = offset_count_type;
+
+	if (expression_type(op2).basetype != offset_count_type)
+	{
+		// Value-cast here. Input might be 16-bit. GLSL requires int.
+		op2_expr = join(type_to_glsl_constructor(target_type), "(", op2_expr, ")");
+	}
+
+	if (expression_type(op3).basetype != offset_count_type)
+	{
+		// Value-cast here. Input might be 16-bit. GLSL requires int.
+		op3_expr = join(type_to_glsl_constructor(target_type), "(", op3_expr, ")");
+	}
+
+	emit_op(result_type, result_id,
+	        join(op, "(", op0_expr, ", ", op1_expr, ", ", op2_expr, ", ", op3_expr, ")"),
+	        forward);
+
+	inherit_expression_dependencies(result_id, op0);
+	inherit_expression_dependencies(result_id, op1);
+	inherit_expression_dependencies(result_id, op2);
+	inherit_expression_dependencies(result_id, op3);
+}
+
 // EXT_shader_texture_lod only concerns fragment shaders so lod tex functions
 // are not allowed in ES 2 vertex shaders. But SPIR-V only supports lod tex
 // functions in vertex shaders so we revert those back to plain calls when
@@ -8981,23 +9072,39 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 
 	// Bitfield
 	case OpBitFieldInsert:
-		// TODO: The signedness of inputs is strict in GLSL, but not in SPIR-V, bitcast if necessary.
-		GLSL_QFOP(bitfieldInsert);
+	{
+		emit_bitfield_insert_op(ops[0], ops[1], ops[2], ops[3], ops[4], ops[5],
+		                        "bitfieldInsert", SPIRType::Int);
 		break;
+	}
 
 	case OpBitFieldSExtract:
-	case OpBitFieldUExtract:
-		// TODO: The signedness of inputs is strict in GLSL, but not in SPIR-V, bitcast if necessary.
-		GLSL_TFOP(bitfieldExtract);
+	{
+		emit_trinary_func_op_bitextract(ops[0], ops[1], ops[2], ops[3], ops[4], "bitfieldExtract",
+		                                int_type, int_type,
+		                                SPIRType::Int, SPIRType::Int);
 		break;
+	}
+
+	case OpBitFieldUExtract:
+	{
+		emit_trinary_func_op_bitextract(ops[0], ops[1], ops[2], ops[3], ops[4], "bitfieldExtract",
+		                                uint_type, uint_type,
+		                                SPIRType::Int, SPIRType::Int);
+		break;
+	}
 
 	case OpBitReverse:
+		// BitReverse does not have issues with sign since result type must match input type.
 		GLSL_UFOP(bitfieldReverse);
 		break;
 
 	case OpBitCount:
-		GLSL_UFOP(bitCount);
+	{
+		auto basetype = expression_type(ops[2]).basetype;
+		emit_unary_func_op_cast(ops[0], ops[1], ops[2], "bitCount", basetype, int_type);
 		break;
+	}
 
 	// Atomics
 	case OpAtomicExchange:

spirv_glsl.hpp

@@ -458,6 +458,16 @@ protected:
 	                              SPIRType::BaseType input_type, bool skip_cast_if_equal_type);
 	void emit_trinary_func_op_cast(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, uint32_t op2,
 	                               const char *op, SPIRType::BaseType input_type);
+	void emit_trinary_func_op_bitextract(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, uint32_t op2,
+	                                     const char *op,
+	                                     SPIRType::BaseType expected_result_type,
+	                                     SPIRType::BaseType input_type0,
+	                                     SPIRType::BaseType input_type1,
+	                                     SPIRType::BaseType input_type2);
+	void emit_bitfield_insert_op(uint32_t result_type, uint32_t result_id,
+	                             uint32_t op0, uint32_t op1, uint32_t op2, uint32_t op3,
+	                             const char *op,
+	                             SPIRType::BaseType offset_count_type);
 
 	void emit_unary_func_op(uint32_t result_type, uint32_t result_id, uint32_t op0, const char *op);
 	void emit_unrolled_unary_op(uint32_t result_type, uint32_t result_id, uint32_t operand, const char *op);

spirv_hlsl.cpp

@@ -4641,8 +4641,11 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
 	}
 
 	case OpBitCount:
-		HLSL_UFOP(countbits);
+	{
+		auto basetype = expression_type(ops[2]).basetype;
+		emit_unary_func_op_cast(ops[0], ops[1], ops[2], "countbits", basetype, basetype);
 		break;
+	}
 
 	case OpBitReverse:
 		HLSL_UFOP(reversebits);

spirv_msl.cpp

@@ -4229,21 +4229,39 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
 
 	// Bitfield
 	case OpBitFieldInsert:
-		MSL_QFOP(insert_bits);
+	{
+		emit_bitfield_insert_op(ops[0], ops[1], ops[2], ops[3], ops[4], ops[5],
+		                        "insert_bits", SPIRType::UInt);
 		break;
+	}
 
 	case OpBitFieldSExtract:
-	case OpBitFieldUExtract:
-		MSL_TFOP(extract_bits);
+	{
+		emit_trinary_func_op_bitextract(ops[0], ops[1], ops[2], ops[3], ops[4], "extract_bits",
+		                                int_type, int_type,
+		                                SPIRType::UInt, SPIRType::UInt);
 		break;
+	}
+
+	case OpBitFieldUExtract:
+	{
+		emit_trinary_func_op_bitextract(ops[0], ops[1], ops[2], ops[3], ops[4], "extract_bits",
+		                                uint_type, uint_type,
+		                                SPIRType::UInt, SPIRType::UInt);
+		break;
+	}
 
 	case OpBitReverse:
+		// BitReverse does not have issues with sign since result type must match input type.
 		MSL_UFOP(reverse_bits);
 		break;
 
 	case OpBitCount:
-		MSL_UFOP(popcount);
+	{
+		auto basetype = expression_type(ops[2]).basetype;
+		emit_unary_func_op_cast(ops[0], ops[1], ops[2], "popcount", basetype, basetype);
 		break;
+	}
 
 	case OpFRem:
 		MSL_BFOP(fmod);