Merge pull request #1724 from billhollings/msl-const-expr-casting

MSL: Fix casting in constant expressions with different sizes.

reference/opt/shaders-msl/comp/type_casting_i64.msl22.comp

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct dst_buff_t
+{
+    int m0[1];
+};
+
+struct src_buff_t
+{
+    int m0[1];
+};
+
+constant int base_val_tmp [[function_constant(0)]];
+constant int base_val = is_function_constant_defined(base_val_tmp) ? base_val_tmp : 0;
+constant long shift_val_tmp [[function_constant(1)]];
+constant long shift_val = is_function_constant_defined(shift_val_tmp) ? shift_val_tmp : 0l;
+constant int offset = (base_val >> int(shift_val));
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
+
+kernel void main0(device dst_buff_t& dst_buff [[buffer(0)]], device src_buff_t& src_buff [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    dst_buff.m0[gl_GlobalInvocationID.x] = src_buff.m0[gl_GlobalInvocationID.x] + offset;
+}
+

reference/shaders-msl/comp/type_casting_i64.msl22.comp

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct dst_buff_t
+{
+    int m0[1];
+};
+
+struct src_buff_t
+{
+    int m0[1];
+};
+
+constant int base_val_tmp [[function_constant(0)]];
+constant int base_val = is_function_constant_defined(base_val_tmp) ? base_val_tmp : 0;
+constant long shift_val_tmp [[function_constant(1)]];
+constant long shift_val = is_function_constant_defined(shift_val_tmp) ? shift_val_tmp : 0l;
+constant int offset = (base_val >> int(shift_val));
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
+
+kernel void main0(device dst_buff_t& dst_buff [[buffer(0)]], device src_buff_t& src_buff [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    dst_buff.m0[gl_GlobalInvocationID.x] = src_buff.m0[gl_GlobalInvocationID.x] + offset;
+}
+

shaders-msl/comp/type_casting_i64.msl22.comp

@@ -0,0 +1,23 @@
+#version 450
+#extension GL_ARB_gpu_shader_int64 : require
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(constant_id = 0) const int base_val = 0;
+layout(constant_id = 1) const int64_t shift_val = 0;
+const int offset = base_val >> shift_val;
+
+layout(set = 0, binding = 0, std430) buffer src_buff_t
+{
+    int m0[];
+} src_buff;
+
+layout(set = 0, binding = 1, std430) buffer dst_buff_t
+{
+    int m0[];
+} dst_buff;
+
+void main()
+{
+    dst_buff.m0[gl_GlobalInvocationID.x] = src_buff.m0[gl_GlobalInvocationID.x] + offset;
+}
+

spirv_msl.cpp

@@ -13880,18 +13880,21 @@ string CompilerMSL::bitcast_glsl_op(const SPIRType &out_type, const SPIRType &in
 	assert(out_type.basetype != SPIRType::Boolean);
 	assert(in_type.basetype != SPIRType::Boolean);
 
-	bool integral_cast = type_is_integral(out_type) && type_is_integral(in_type);
-	bool same_size_cast = out_type.width == in_type.width;
+	bool integral_cast = type_is_integral(out_type) && type_is_integral(in_type) && (out_type.vecsize == in_type.vecsize);
+	bool same_size_cast = (out_type.width * out_type.vecsize) == (in_type.width * in_type.vecsize);
 
-	if (integral_cast && same_size_cast)
+	// Bitcasting can only be used between types of the same overall size.
+	// And always formally cast between integers, because it's trivial, and also
+	// because Metal can internally cast the results of some integer ops to a larger
+	// size (eg. short shift right becomes int), which means chaining integer ops
+	// together may introduce size variations that SPIR-V doesn't know about.
+	if (same_size_cast && !integral_cast)
 	{
-		// Trivial bitcast case, casts between integers.
-		return type_to_glsl(out_type);
+		return "as_type<" + type_to_glsl(out_type) + ">";
 	}
 	else
 	{
-		// Fall back to the catch-all bitcast in MSL.
-		return "as_type<" + type_to_glsl(out_type) + ">";
+		return type_to_glsl(out_type);
 	}
 }