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 <metal_stdlib>
+#include <simd/simd.h>
+
+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<ulong4>(as_type<double4>(ssbo_0.a));
+    ssbo_0.a = as_type<long4>(as_type<double4>(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 <metal_stdlib>
+#include <simd/simd.h>
+
+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<ulong4>(as_type<double4>(ssbo_0.a));
+    ssbo_0.a = as_type<long4>(as_type<double4>(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";