From a75fe07546c2d5d584208d76ae84147b413a3fb0 Mon Sep 17 00:00:00 2001 From: Bill Hollings Date: Thu, 12 Aug 2021 16:19:46 -0400 Subject: [PATCH 1/2] MSL: Fix casting in constant expressions with different sizes. Previous casting in constant expressions used as_type<> between types of different overall sizes. Add check for overall size (width * vecsize) to ensure as_type<> will work, otherwise use regular cast. Also beef up test of integer values to also check vecsize, and use regular casts for those. --- spirv_msl.cpp | 17 ++++++++++------- 1 file changed, 10 insertions(+), 7 deletions(-) diff --git a/spirv_msl.cpp b/spirv_msl.cpp index 76c55652..9eef7c5d 100644 --- a/spirv_msl.cpp +++ b/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); } } From e76fcf93099334ed47277b4427e89fc2b8956241 Mon Sep 17 00:00:00 2001 From: Bill Hollings Date: Mon, 16 Aug 2021 13:56:05 -0400 Subject: [PATCH 2/2] MSL: Add test for fixes to MSL constant expression type down-casting. --- .../comp/type_casting_i64.msl22.comp | 27 +++++++++++++++++++ .../comp/type_casting_i64.msl22.comp | 27 +++++++++++++++++++ shaders-msl/comp/type_casting_i64.msl22.comp | 23 ++++++++++++++++ 3 files changed, 77 insertions(+) create mode 100644 reference/opt/shaders-msl/comp/type_casting_i64.msl22.comp create mode 100644 reference/shaders-msl/comp/type_casting_i64.msl22.comp create mode 100644 shaders-msl/comp/type_casting_i64.msl22.comp diff --git a/reference/opt/shaders-msl/comp/type_casting_i64.msl22.comp b/reference/opt/shaders-msl/comp/type_casting_i64.msl22.comp new file mode 100644 index 00000000..6820b077 --- /dev/null +++ b/reference/opt/shaders-msl/comp/type_casting_i64.msl22.comp @@ -0,0 +1,27 @@ +#include +#include + +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; +} + diff --git a/reference/shaders-msl/comp/type_casting_i64.msl22.comp b/reference/shaders-msl/comp/type_casting_i64.msl22.comp new file mode 100644 index 00000000..6820b077 --- /dev/null +++ b/reference/shaders-msl/comp/type_casting_i64.msl22.comp @@ -0,0 +1,27 @@ +#include +#include + +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; +} + diff --git a/shaders-msl/comp/type_casting_i64.msl22.comp b/shaders-msl/comp/type_casting_i64.msl22.comp new file mode 100644 index 00000000..45e682e5 --- /dev/null +++ b/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; +} +