AuroraMiddleware/SPIRV-Cross
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Merge pull request #1031 from KhronosGroup/fix-1009

Browse Source 
MSL: Support 64-bit integers.
This commit is contained in:
Hans-Kristian Arntzen 2019-06-19 15:29:27 +02:00 committed by GitHub
commit a1f7c8dc8e
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
4 changed files with 165 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

24
reference/opt/shaders-msl/comp/int64.invalid.msl22.comp Normal file
View File

@ -0,0 +1,24 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct SSBO
{
int s32;
uint u32;
};
constant long _162 = {};
kernel void main0(device SSBO& _96 [[buffer(0)]])
{
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);
}

63
reference/shaders-msl/comp/int64.invalid.msl22.comp Normal file
View File

@ -0,0 +1,63 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct M0
{
long v;
long2 b[2];
ulong c;
ulong d[5];
};
struct SSBO0_Type
{
long4 a;
M0 m0;
};
struct SSBO1_Type
{
ulong4 b;
M0 m0;
};
struct SSBO2_Type
{
long a[4];
long2 b[4];
};
struct SSBO3_Type
{
long a[4];
long2 b[4];
};
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));
ssbo_0.a += long4(1l);
ssbo_1.b += ulong4(long4(1l));
ssbo_0.a -= long4(1l);
ssbo_1.b -= ulong4(long4(1l));
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]));
}

65
shaders-msl/comp/int64.invalid.msl22.comp Normal file
View File

@ -0,0 +1,65 @@
#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];
};
struct SSBO0_Type
{
i64vec4 a;
M0 m0;
};
struct SSBO1_Type
{
u64vec4 b;
M0 m0;
};
struct SSBO2_Type
{
int64_t a[4];
i64vec2 b[4];
};
struct SSBO3_Type
{
int64_t a[4];
i64vec2 b[4];
};
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;
ssbo_0.a = abs(ssbo_0.a + i64vec4(ssbo_1.b));
ssbo_0.a++;
ssbo_1.b++;
ssbo_0.a--;
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]);
}

15
spirv_msl.cpp
View File

@ -7094,10 +7094,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";
@ -8005,6 +8009,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.