From a6b71ae9990cb0a1b17b1814e19ccd3a04b64907 Mon Sep 17 00:00:00 2001 From: Hans-Kristian Arntzen Date: Tue, 11 Jun 2019 10:45:22 +0200 Subject: [PATCH 1/2] MSL: Support 64-bit integers. --- .../shaders-msl/comp/int64.invalid.msl22.comp | 53 ++++++++++++++++++ .../shaders-msl/comp/int64.invalid.msl22.comp | 53 ++++++++++++++++++ shaders-msl/comp/int64.invalid.msl22.comp | 55 +++++++++++++++++++ spirv_msl.cpp | 8 ++- 4 files changed, 167 insertions(+), 2 deletions(-) create mode 100644 reference/opt/shaders-msl/comp/int64.invalid.msl22.comp create mode 100644 reference/shaders-msl/comp/int64.invalid.msl22.comp create mode 100644 shaders-msl/comp/int64.invalid.msl22.comp diff --git a/reference/opt/shaders-msl/comp/int64.invalid.msl22.comp b/reference/opt/shaders-msl/comp/int64.invalid.msl22.comp new file mode 100644 index 00000000..7e69ddc3 --- /dev/null +++ b/reference/opt/shaders-msl/comp/int64.invalid.msl22.comp @@ -0,0 +1,53 @@ +#include +#include + +using namespace metal; + +struct M0 +{ + long v; + long2 b[2]; + ulong c; + ulong d[5]; +}; + +struct SSBO0 +{ + long4 a; + M0 m0; +}; + +struct SSBO1 +{ + ulong4 b; + M0 m0; +}; + +struct SSBO2 +{ + long a[4]; + long2 b[4]; +}; + +struct SSBO3 +{ + long a[4]; + long2 b[4]; +}; + +kernel void main0(device SSBO0& ssbo_0 [[buffer(0)]], device SSBO1& ssbo_1 [[buffer(1)]], device SSBO2& ssbo_2 [[buffer(2)]], device SSBO3& ssbo_3 [[buffer(3)]]) +{ + ssbo_0.a += long4(10l, 20l, 30l, 40l); + ssbo_1.b += ulong4(999999999999999999ul, 8888888888888888ul, 77777777777777777ul, 6666666666666666ul); + ssbo_0.a += long4(20l); + ssbo_0.a = abs(ssbo_0.a + long4(ssbo_1.b)); + ssbo_0.a += long4(1l); + ssbo_1.b += ulong4(long4(1l)); + ssbo_0.a -= long4(1l); + ssbo_1.b -= ulong4(long4(1l)); + ssbo_1.b = as_type(as_type(ssbo_0.a)); + ssbo_0.a = as_type(as_type(ssbo_1.b)); + ssbo_2.a[0] += 1l; + ssbo_3.a[0] += 2l; +} + diff --git a/reference/shaders-msl/comp/int64.invalid.msl22.comp b/reference/shaders-msl/comp/int64.invalid.msl22.comp new file mode 100644 index 00000000..7e69ddc3 --- /dev/null +++ b/reference/shaders-msl/comp/int64.invalid.msl22.comp @@ -0,0 +1,53 @@ +#include +#include + +using namespace metal; + +struct M0 +{ + long v; + long2 b[2]; + ulong c; + ulong d[5]; +}; + +struct SSBO0 +{ + long4 a; + M0 m0; +}; + +struct SSBO1 +{ + ulong4 b; + M0 m0; +}; + +struct SSBO2 +{ + long a[4]; + long2 b[4]; +}; + +struct SSBO3 +{ + long a[4]; + long2 b[4]; +}; + +kernel void main0(device SSBO0& ssbo_0 [[buffer(0)]], device SSBO1& ssbo_1 [[buffer(1)]], device SSBO2& ssbo_2 [[buffer(2)]], device SSBO3& ssbo_3 [[buffer(3)]]) +{ + ssbo_0.a += long4(10l, 20l, 30l, 40l); + ssbo_1.b += ulong4(999999999999999999ul, 8888888888888888ul, 77777777777777777ul, 6666666666666666ul); + ssbo_0.a += long4(20l); + ssbo_0.a = abs(ssbo_0.a + long4(ssbo_1.b)); + ssbo_0.a += long4(1l); + ssbo_1.b += ulong4(long4(1l)); + ssbo_0.a -= long4(1l); + ssbo_1.b -= ulong4(long4(1l)); + ssbo_1.b = as_type(as_type(ssbo_0.a)); + ssbo_0.a = as_type(as_type(ssbo_1.b)); + ssbo_2.a[0] += 1l; + ssbo_3.a[0] += 2l; +} + diff --git a/shaders-msl/comp/int64.invalid.msl22.comp b/shaders-msl/comp/int64.invalid.msl22.comp new file mode 100644 index 00000000..81004d4a --- /dev/null +++ b/shaders-msl/comp/int64.invalid.msl22.comp @@ -0,0 +1,55 @@ +#version 450 +#extension GL_ARB_gpu_shader_int64 : require +layout(local_size_x = 1) in; + +struct M0 +{ + int64_t v; + i64vec2 b[2]; + uint64_t c; + uint64_t d[5]; +}; + +// Test buffer layout handling. +layout(std430, binding = 0) buffer SSBO0 +{ + i64vec4 a; + M0 m0; +} ssbo_0; + +layout(std430, binding = 1) buffer SSBO1 +{ + u64vec4 b; + M0 m0; +} ssbo_1; + +layout(std430, binding = 2) buffer SSBO2 +{ + int64_t a[4]; + i64vec2 b[4]; +} ssbo_2; + +layout(std140, binding = 3) buffer SSBO3 +{ + int64_t a[4]; + i64vec2 b[4]; +} ssbo_3; + +void main() +{ + ssbo_0.a += i64vec4(10, 20, 30, 40); + ssbo_1.b += u64vec4(999999999999999999ul, 8888888888888888ul, 77777777777777777ul, 6666666666666666ul); + ssbo_0.a += 20; + ssbo_0.a = abs(ssbo_0.a + i64vec4(ssbo_1.b)); + + ssbo_0.a++; + ssbo_1.b++; + ssbo_0.a--; + ssbo_1.b--; + + ssbo_1.b = doubleBitsToUint64(int64BitsToDouble(ssbo_0.a)); + ssbo_0.a = doubleBitsToInt64(uint64BitsToDouble(ssbo_1.b)); + + ssbo_2.a[0] += 1l; + ssbo_3.a[0] += 2l; +} diff --git a/spirv_msl.cpp b/spirv_msl.cpp index 3a827c65..130234c9 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -7024,10 +7024,14 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id) type_name = "uint"; break; case SPIRType::Int64: - type_name = "long"; // Currently unsupported + if (!msl_options.supports_msl_version(2, 2)) + SPIRV_CROSS_THROW("64-bit integers are only supported in MSL 2.2 and above."); + type_name = "long"; break; case SPIRType::UInt64: - type_name = "size_t"; + if (!msl_options.supports_msl_version(2, 2)) + SPIRV_CROSS_THROW("64-bit integers are only supported in MSL 2.2 and above."); + type_name = "ulong"; break; case SPIRType::Half: type_name = "half"; From a6798d06a2866c5d4f87122d61a542e2857ac6ff Mon Sep 17 00:00:00 2001 From: Hans-Kristian Arntzen Date: Wed, 19 Jun 2019 10:11:10 +0200 Subject: [PATCH 2/2] MSL: Error out on int64_t/uint64_t buffer members. Not supported for whatever reason. --- .../shaders-msl/comp/int64.invalid.msl22.comp | 53 +++++-------------- .../shaders-msl/comp/int64.invalid.msl22.comp | 24 ++++++--- shaders-msl/comp/int64.invalid.msl22.comp | 34 +++++++----- spirv_msl.cpp | 7 +++ 4 files changed, 58 insertions(+), 60 deletions(-) diff --git a/reference/opt/shaders-msl/comp/int64.invalid.msl22.comp b/reference/opt/shaders-msl/comp/int64.invalid.msl22.comp index 7e69ddc3..13304bd0 100644 --- a/reference/opt/shaders-msl/comp/int64.invalid.msl22.comp +++ b/reference/opt/shaders-msl/comp/int64.invalid.msl22.comp @@ -3,51 +3,22 @@ using namespace metal; -struct M0 +struct SSBO { - long v; - long2 b[2]; - ulong c; - ulong d[5]; + int s32; + uint u32; }; -struct SSBO0 -{ - long4 a; - M0 m0; -}; +constant long _162 = {}; -struct SSBO1 +kernel void main0(device SSBO& _96 [[buffer(0)]]) { - ulong4 b; - M0 m0; -}; - -struct SSBO2 -{ - long a[4]; - long2 b[4]; -}; - -struct SSBO3 -{ - long a[4]; - long2 b[4]; -}; - -kernel void main0(device SSBO0& ssbo_0 [[buffer(0)]], device SSBO1& ssbo_1 [[buffer(1)]], device SSBO2& ssbo_2 [[buffer(2)]], device SSBO3& ssbo_3 [[buffer(3)]]) -{ - ssbo_0.a += long4(10l, 20l, 30l, 40l); - ssbo_1.b += ulong4(999999999999999999ul, 8888888888888888ul, 77777777777777777ul, 6666666666666666ul); - ssbo_0.a += long4(20l); - ssbo_0.a = abs(ssbo_0.a + long4(ssbo_1.b)); - ssbo_0.a += long4(1l); - ssbo_1.b += ulong4(long4(1l)); - ssbo_0.a -= long4(1l); - ssbo_1.b -= ulong4(long4(1l)); - ssbo_1.b = as_type(as_type(ssbo_0.a)); - ssbo_0.a = as_type(as_type(ssbo_1.b)); - ssbo_2.a[0] += 1l; - ssbo_3.a[0] += 2l; + long4 _137; + ulong4 _141; + _137 = abs((_137 + long4(30l, 40l, 50l, 60l)) + long4(_141 + ulong4(999999999999999999ul, 8888888888888888ul, 77777777777777777ul, 6666666666666666ul))); + _141 += ulong4(long4(999999999999999999l, 8888888888888888l, 77777777777777777l, 6666666666666666l)); + ulong _109 = ulong(_162); + _96.s32 = int(uint(((ulong(_137.x) + _141.y) + _109) + _109)); + _96.u32 = uint(((ulong(_137.y) + _141.z) + ulong(_162 + 1l)) + _109); } diff --git a/reference/shaders-msl/comp/int64.invalid.msl22.comp b/reference/shaders-msl/comp/int64.invalid.msl22.comp index 7e69ddc3..6eb4a8a8 100644 --- a/reference/shaders-msl/comp/int64.invalid.msl22.comp +++ b/reference/shaders-msl/comp/int64.invalid.msl22.comp @@ -11,33 +11,41 @@ struct M0 ulong d[5]; }; -struct SSBO0 +struct SSBO0_Type { long4 a; M0 m0; }; -struct SSBO1 +struct SSBO1_Type { ulong4 b; M0 m0; }; -struct SSBO2 +struct SSBO2_Type { long a[4]; long2 b[4]; }; -struct SSBO3 +struct SSBO3_Type { long a[4]; long2 b[4]; }; -kernel void main0(device SSBO0& ssbo_0 [[buffer(0)]], device SSBO1& ssbo_1 [[buffer(1)]], device SSBO2& ssbo_2 [[buffer(2)]], device SSBO3& ssbo_3 [[buffer(3)]]) +struct SSBO { + int s32; + uint u32; +}; + +kernel void main0(device SSBO& _96 [[buffer(0)]]) +{ + SSBO0_Type ssbo_0; ssbo_0.a += long4(10l, 20l, 30l, 40l); + SSBO1_Type ssbo_1; ssbo_1.b += ulong4(999999999999999999ul, 8888888888888888ul, 77777777777777777ul, 6666666666666666ul); ssbo_0.a += long4(20l); ssbo_0.a = abs(ssbo_0.a + long4(ssbo_1.b)); @@ -45,9 +53,11 @@ kernel void main0(device SSBO0& ssbo_0 [[buffer(0)]], device SSBO1& ssbo_1 [[buf ssbo_1.b += ulong4(long4(1l)); ssbo_0.a -= long4(1l); ssbo_1.b -= ulong4(long4(1l)); - ssbo_1.b = as_type(as_type(ssbo_0.a)); - ssbo_0.a = as_type(as_type(ssbo_1.b)); + SSBO2_Type ssbo_2; ssbo_2.a[0] += 1l; + SSBO3_Type ssbo_3; ssbo_3.a[0] += 2l; + _96.s32 = int(uint(((ulong(ssbo_0.a.x) + ssbo_1.b.y) + ulong(ssbo_2.a[1])) + ulong(ssbo_3.a[2]))); + _96.u32 = uint(((ulong(ssbo_0.a.y) + ssbo_1.b.z) + ulong(ssbo_2.a[0])) + ulong(ssbo_3.a[1])); } diff --git a/shaders-msl/comp/int64.invalid.msl22.comp b/shaders-msl/comp/int64.invalid.msl22.comp index 81004d4a..965bed4a 100644 --- a/shaders-msl/comp/int64.invalid.msl22.comp +++ b/shaders-msl/comp/int64.invalid.msl22.comp @@ -10,33 +10,43 @@ struct M0 uint64_t d[5]; }; -// Test buffer layout handling. -layout(std430, binding = 0) buffer SSBO0 +struct SSBO0_Type { i64vec4 a; M0 m0; -} ssbo_0; +}; -layout(std430, binding = 1) buffer SSBO1 +struct SSBO1_Type { u64vec4 b; M0 m0; -} ssbo_1; +}; -layout(std430, binding = 2) buffer SSBO2 +struct SSBO2_Type { int64_t a[4]; i64vec2 b[4]; -} ssbo_2; +}; -layout(std140, binding = 3) buffer SSBO3 +struct SSBO3_Type { int64_t a[4]; i64vec2 b[4]; -} ssbo_3; +}; + +layout(set = 0, binding = 0, std430) buffer SSBO +{ + int s32; + uint u32; +}; void main() { + SSBO0_Type ssbo_0; + SSBO1_Type ssbo_1; + SSBO2_Type ssbo_2; + SSBO3_Type ssbo_3; + ssbo_0.a += i64vec4(10, 20, 30, 40); ssbo_1.b += u64vec4(999999999999999999ul, 8888888888888888ul, 77777777777777777ul, 6666666666666666ul); ssbo_0.a += 20; @@ -47,9 +57,9 @@ void main() ssbo_0.a--; ssbo_1.b--; - ssbo_1.b = doubleBitsToUint64(int64BitsToDouble(ssbo_0.a)); - ssbo_0.a = doubleBitsToInt64(uint64BitsToDouble(ssbo_1.b)); - ssbo_2.a[0] += 1l; ssbo_3.a[0] += 2l; + + s32 = int(ssbo_0.a.x + ssbo_1.b.y + ssbo_2.a[1] + ssbo_3.a[2]); + u32 = uint(ssbo_0.a.y + ssbo_1.b.z + ssbo_2.a[0] + ssbo_3.a[1]); } diff --git a/spirv_msl.cpp b/spirv_msl.cpp index 130234c9..495a91d4 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -7906,6 +7906,13 @@ size_t CompilerMSL::get_declared_struct_member_alignment(const SPIRType &struct_ case SPIRType::Sampler: SPIRV_CROSS_THROW("Querying alignment of opaque object."); + case SPIRType::Int64: + SPIRV_CROSS_THROW("long types are not supported in buffers in MSL."); + case SPIRType::UInt64: + SPIRV_CROSS_THROW("ulong types are not supported in buffers in MSL."); + case SPIRType::Double: + SPIRV_CROSS_THROW("double types are not supported in buffers in MSL."); + case SPIRType::Struct: { // In MSL, a struct's alignment is equal to the maximum alignment of any of its members.