Declare read-only SSBOs as const device in MSL.

reference/opt/shaders-msl/comp/basic.comp

@@ -21,7 +21,7 @@ struct SSBO3
     uint counter;
 };
 
-kernel void main0(device SSBO& _23 [[buffer(0)]], device SSBO2& _45 [[buffer(1)]], device SSBO3& _48 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _45 [[buffer(1)]], device SSBO3& _48 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     float4 _29 = _23.in_data[gl_GlobalInvocationID.x];
     if (dot(_29, float4(1.0, 5.0, 6.0, 2.0)) > 8.19999980926513671875)

reference/opt/shaders-msl/comp/culling.comp

@@ -23,7 +23,7 @@ struct SSBO3
     uint count;
 };
 
-kernel void main0(device SSBO& _22 [[buffer(0)]], device SSBO2& _38 [[buffer(1)]], device SSBO3& _41 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(const device SSBO& _22 [[buffer(0)]], device SSBO2& _38 [[buffer(1)]], device SSBO3& _41 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     float _28 = _22.in_data[gl_GlobalInvocationID.x];
     if (_28 > 12.0)

reference/opt/shaders-msl/comp/dowhile.comp

@@ -14,7 +14,7 @@ struct SSBO2
     float4 out_data[1];
 };
 
-kernel void main0(device SSBO& _28 [[buffer(0)]], device SSBO2& _52 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(const device SSBO& _28 [[buffer(0)]], device SSBO2& _52 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     int i = 0;
     float4 _56;

reference/opt/shaders-msl/comp/inverse.comp

@@ -114,7 +114,7 @@ float2x2 spvInverse2x2(float2x2 m)
     return (det != 0.0f) ? (adj * (1.0f / det)) : m;
 }
 
-kernel void main0(device MatrixOut& _15 [[buffer(0)]], device MatrixIn& _20 [[buffer(1)]])
+kernel void main0(device MatrixOut& _15 [[buffer(0)]], const device MatrixIn& _20 [[buffer(1)]])
 {
     _15.m2out = spvInverse2x2(_20.m2in);
     _15.m3out = spvInverse3x3(_20.m3in);

reference/opt/shaders-msl/comp/mod.comp

@@ -22,7 +22,7 @@ Tx mod(Tx x, Ty y)
     return x - y * floor(x / y);
 }
 
-kernel void main0(device SSBO& _23 [[buffer(0)]], device SSBO2& _33 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _33 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     _33.out_data[gl_GlobalInvocationID.x] = mod(_23.in_data[gl_GlobalInvocationID.x], _33.out_data[gl_GlobalInvocationID.x]);
     _33.out_data[gl_GlobalInvocationID.x] = as_type<float4>(as_type<uint4>(_23.in_data[gl_GlobalInvocationID.x]) % as_type<uint4>(_33.out_data[gl_GlobalInvocationID.x]));

reference/opt/shaders-msl/comp/modf.comp

@@ -13,7 +13,7 @@ struct SSBO2
     float4 out_data[1];
 };
 
-kernel void main0(device SSBO& _23 [[buffer(0)]], device SSBO2& _35 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _35 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     float4 i;
     float4 _31 = modf(_23.in_data[gl_GlobalInvocationID.x], i);

reference/opt/shaders-msl/comp/read-write-only.comp

@@ -21,7 +21,7 @@ struct SSBO1
     float4 data3;
 };
 
-kernel void main0(device SSBO0& _15 [[buffer(0)]], device SSBO1& _21 [[buffer(1)]], device SSBO2& _10 [[buffer(2)]])
+kernel void main0(const device SSBO0& _15 [[buffer(0)]], device SSBO1& _21 [[buffer(1)]], device SSBO2& _10 [[buffer(2)]])
 {
     _10.data4 = _15.data0 + _21.data2;
     _10.data5 = _15.data1 + _21.data3;

reference/opt/shaders-msl/comp/shared.comp

@@ -15,7 +15,7 @@ struct SSBO2
     float out_data[1];
 };
 
-kernel void main0(device SSBO& _22 [[buffer(0)]], device SSBO2& _44 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
+kernel void main0(const device SSBO& _22 [[buffer(0)]], device SSBO2& _44 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
 {
     threadgroup float sShared[4];
     sShared[gl_LocalInvocationIndex] = _22.in_data[gl_GlobalInvocationID.x];

reference/opt/shaders-msl/comp/struct-layout.comp

@@ -18,7 +18,7 @@ struct SSBO
     Foo in_data[1];
 };
 
-kernel void main0(device SSBO& _30 [[buffer(0)]], device SSBO2& _23 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(const device SSBO& _30 [[buffer(0)]], device SSBO2& _23 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     _23.out_data[gl_GlobalInvocationID.x].m = _30.in_data[gl_GlobalInvocationID.x].m * _30.in_data[gl_GlobalInvocationID.x].m;
 }

reference/opt/shaders-msl/comp/torture-loop.comp

@@ -14,7 +14,7 @@ struct SSBO2
     float4 out_data[1];
 };
 
-kernel void main0(device SSBO& _24 [[buffer(0)]], device SSBO2& _89 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(const device SSBO& _24 [[buffer(0)]], device SSBO2& _89 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     float4 _99;
     _99 = _24.in_data[gl_GlobalInvocationID.x];

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

@@ -14,7 +14,7 @@ struct SSBO2
     float4 out_data[1];
 };
 
-kernel void main0(device SSBO& _24 [[buffer(0)]], device SSBO2& _177 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(const device SSBO& _24 [[buffer(0)]], device SSBO2& _177 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     uint ident = gl_GlobalInvocationID.x;
     float4 idat = _24.in_data[ident];

reference/shaders-msl/comp/basic.comp

@@ -21,7 +21,7 @@ struct SSBO3
     uint counter;
 };
 
-kernel void main0(device SSBO& _23 [[buffer(0)]], device SSBO2& _45 [[buffer(1)]], device SSBO3& _48 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _45 [[buffer(1)]], device SSBO3& _48 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     uint ident = gl_GlobalInvocationID.x;
     float4 idata = _23.in_data[ident];

reference/shaders-msl/comp/culling.comp

@@ -23,7 +23,7 @@ struct SSBO3
     uint count;
 };
 
-kernel void main0(device SSBO& _22 [[buffer(0)]], device SSBO2& _38 [[buffer(1)]], device SSBO3& _41 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(const device SSBO& _22 [[buffer(0)]], device SSBO2& _38 [[buffer(1)]], device SSBO3& _41 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     uint ident = gl_GlobalInvocationID.x;
     float idata = _22.in_data[ident];

reference/shaders-msl/comp/dowhile.comp

@@ -14,7 +14,7 @@ struct SSBO2
     float4 out_data[1];
 };
 
-kernel void main0(device SSBO& _28 [[buffer(0)]], device SSBO2& _52 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(const device SSBO& _28 [[buffer(0)]], device SSBO2& _52 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     uint ident = gl_GlobalInvocationID.x;
     int i = 0;

reference/shaders-msl/comp/inverse.comp

@@ -114,7 +114,7 @@ float2x2 spvInverse2x2(float2x2 m)
     return (det != 0.0f) ? (adj * (1.0f / det)) : m;
 }
 
-kernel void main0(device MatrixOut& _15 [[buffer(0)]], device MatrixIn& _20 [[buffer(1)]])
+kernel void main0(device MatrixOut& _15 [[buffer(0)]], const device MatrixIn& _20 [[buffer(1)]])
 {
     _15.m2out = spvInverse2x2(_20.m2in);
     _15.m3out = spvInverse3x3(_20.m3in);

reference/shaders-msl/comp/mod.comp

@@ -22,7 +22,7 @@ Tx mod(Tx x, Ty y)
     return x - y * floor(x / y);
 }
 
-kernel void main0(device SSBO& _23 [[buffer(0)]], device SSBO2& _33 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _33 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     uint ident = gl_GlobalInvocationID.x;
     float4 v = mod(_23.in_data[ident], _33.out_data[ident]);

reference/shaders-msl/comp/modf.comp

@@ -13,7 +13,7 @@ struct SSBO2
     float4 out_data[1];
 };
 
-kernel void main0(device SSBO& _23 [[buffer(0)]], device SSBO2& _35 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _35 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     uint ident = gl_GlobalInvocationID.x;
     float4 i;

reference/shaders-msl/comp/read-write-only.comp

@@ -21,7 +21,7 @@ struct SSBO1
     float4 data3;
 };
 
-kernel void main0(device SSBO0& _15 [[buffer(0)]], device SSBO1& _21 [[buffer(1)]], device SSBO2& _10 [[buffer(2)]])
+kernel void main0(const device SSBO0& _15 [[buffer(0)]], device SSBO1& _21 [[buffer(1)]], device SSBO2& _10 [[buffer(2)]])
 {
     _10.data4 = _15.data0 + _21.data2;
     _10.data5 = _15.data1 + _21.data3;

reference/shaders-msl/comp/shared.comp

@@ -15,7 +15,7 @@ struct SSBO2
     float out_data[1];
 };
 
-kernel void main0(device SSBO& _22 [[buffer(0)]], device SSBO2& _44 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
+kernel void main0(const device SSBO& _22 [[buffer(0)]], device SSBO2& _44 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
 {
     threadgroup float sShared[4];
     uint ident = gl_GlobalInvocationID.x;

reference/shaders-msl/comp/struct-layout.comp

@@ -18,7 +18,7 @@ struct SSBO
     Foo in_data[1];
 };
 
-kernel void main0(device SSBO& _30 [[buffer(0)]], device SSBO2& _23 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(const device SSBO& _30 [[buffer(0)]], device SSBO2& _23 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     uint ident = gl_GlobalInvocationID.x;
     _23.out_data[ident].m = _30.in_data[ident].m * _30.in_data[ident].m;

reference/shaders-msl/comp/torture-loop.comp

@@ -14,7 +14,7 @@ struct SSBO2
     float4 out_data[1];
 };
 
-kernel void main0(device SSBO& _24 [[buffer(0)]], device SSBO2& _89 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(const device SSBO& _24 [[buffer(0)]], device SSBO2& _89 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     uint ident = gl_GlobalInvocationID.x;
     float4 idat = _24.in_data[ident];

spirv_msl.cpp

@@ -3190,17 +3190,25 @@ string CompilerMSL::get_argument_address_space(const SPIRVariable &argument)
 		return "threadgroup";
 
 	case StorageClassStorageBuffer:
-		return "device";
+	{
+		auto flags = get_buffer_block_flags(argument);
+		return flags.get(DecorationNonWritable) ? "const device" : "device";
+	}
 
 	case StorageClassUniform:
 	case StorageClassUniformConstant:
 	case StorageClassPushConstant:
 		if (type.basetype == SPIRType::Struct)
-			return (meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock) &&
+		{
-			        !meta[argument.self].decoration.decoration_flags.get(DecorationNonWritable)) ?
+			bool ssbo = has_decoration(type.self, DecorationBufferBlock);
-			           "device" :
+			if (!ssbo)
-			           "constant";
+				return "constant";
-
+			else
+			{
+				bool readonly = get_buffer_block_flags(argument).get(DecorationNonWritable);
+				return readonly ? "const device" : "device";
+			}
+		}
 		break;
 
 	default: