diff --git a/reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.comp b/reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.comp
new file mode 100644
index 00000000..a8ade54b
--- /dev/null
+++ b/reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.comp
@@ -0,0 +1,25 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+struct SSBO
+{
+    float4 outdata;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
+
+// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
+#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y)
+
+kernel void main0(device SSBO& _31 [[buffer(1)]], texture2d<uint> uImage [[texture(0)]], device atomic_uint* uImage_atomic [[buffer(0)]], texture2d<float> uTexture [[texture(1)]], sampler uTextureSmplr [[sampler(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    uint _26 = atomic_fetch_add_explicit((device atomic_uint*)&uImage_atomic[spvImage2DAtomicCoord(int2(gl_GlobalInvocationID.xy), uImage)], 10u, memory_order_relaxed);
+    _31.outdata = uTexture.sample(uTextureSmplr, float2(gl_GlobalInvocationID.xy), level(0.0)) + float4(float(_26));
+}
+
diff --git a/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.frag b/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.frag
index a416259d..98cdda44 100644
--- a/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.frag
+++ b/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.frag
@@ -68,7 +68,7 @@ struct main0_in
     uint in_var_TEXCOORD0 [[user(locn0)]];
 };
 
-fragment main0_out main0(main0_in in [[stage_in]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d<uint> RWShadowTileNumCulledObjects [[texture(2)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]])
+fragment main0_out main0(main0_in in [[stage_in]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d<uint> RWShadowTileNumCulledObjects [[texture(0)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]])
 {
     main0_out out = {};
     uint2 _77 = uint2(gl_FragCoord.xy);
diff --git a/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.graphics-robust-access.frag b/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.graphics-robust-access.frag
index 85e27bd1..79558139 100644
--- a/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.graphics-robust-access.frag
+++ b/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.graphics-robust-access.frag
@@ -68,7 +68,7 @@ struct main0_in
     uint in_var_TEXCOORD0 [[user(locn0)]];
 };
 
-fragment main0_out main0(main0_in in [[stage_in]], constant uint* spvBufferSizeConstants [[buffer(25)]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d<uint> RWShadowTileNumCulledObjects [[texture(2)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]])
+fragment main0_out main0(main0_in in [[stage_in]], constant uint* spvBufferSizeConstants [[buffer(25)]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d<uint> RWShadowTileNumCulledObjects [[texture(0)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]])
 {
     main0_out out = {};
     constant uint& CulledObjectBoxBoundsBufferSize = spvBufferSizeConstants[0];
diff --git a/reference/shaders-msl/comp/image-atomic-automatic-bindings.comp b/reference/shaders-msl/comp/image-atomic-automatic-bindings.comp
new file mode 100644
index 00000000..40a8dfcd
--- /dev/null
+++ b/reference/shaders-msl/comp/image-atomic-automatic-bindings.comp
@@ -0,0 +1,26 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+struct SSBO
+{
+    float4 outdata;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
+
+// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
+#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y)
+
+kernel void main0(device SSBO& _31 [[buffer(1)]], texture2d<uint> uImage [[texture(0)]], device atomic_uint* uImage_atomic [[buffer(0)]], texture2d<float> uTexture [[texture(1)]], sampler uTextureSmplr [[sampler(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    uint _26 = atomic_fetch_add_explicit((device atomic_uint*)&uImage_atomic[spvImage2DAtomicCoord(int2(gl_GlobalInvocationID.xy), uImage)], 10u, memory_order_relaxed);
+    uint ret = _26;
+    _31.outdata = uTexture.sample(uTextureSmplr, float2(gl_GlobalInvocationID.xy), level(0.0)) + float4(float(ret));
+}
+
diff --git a/reference/shaders-ue4/asm/frag/texture-atomics.asm.frag b/reference/shaders-ue4/asm/frag/texture-atomics.asm.frag
index a416259d..98cdda44 100644
--- a/reference/shaders-ue4/asm/frag/texture-atomics.asm.frag
+++ b/reference/shaders-ue4/asm/frag/texture-atomics.asm.frag
@@ -68,7 +68,7 @@ struct main0_in
     uint in_var_TEXCOORD0 [[user(locn0)]];
 };
 
-fragment main0_out main0(main0_in in [[stage_in]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d<uint> RWShadowTileNumCulledObjects [[texture(2)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]])
+fragment main0_out main0(main0_in in [[stage_in]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d<uint> RWShadowTileNumCulledObjects [[texture(0)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]])
 {
     main0_out out = {};
     uint2 _77 = uint2(gl_FragCoord.xy);
diff --git a/reference/shaders-ue4/asm/frag/texture-atomics.asm.graphics-robust-access.frag b/reference/shaders-ue4/asm/frag/texture-atomics.asm.graphics-robust-access.frag
index a416259d..98cdda44 100644
--- a/reference/shaders-ue4/asm/frag/texture-atomics.asm.graphics-robust-access.frag
+++ b/reference/shaders-ue4/asm/frag/texture-atomics.asm.graphics-robust-access.frag
@@ -68,7 +68,7 @@ struct main0_in
     uint in_var_TEXCOORD0 [[user(locn0)]];
 };
 
-fragment main0_out main0(main0_in in [[stage_in]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d<uint> RWShadowTileNumCulledObjects [[texture(2)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]])
+fragment main0_out main0(main0_in in [[stage_in]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d<uint> RWShadowTileNumCulledObjects [[texture(0)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]])
 {
     main0_out out = {};
     uint2 _77 = uint2(gl_FragCoord.xy);
diff --git a/shaders-msl/comp/image-atomic-automatic-bindings.comp b/shaders-msl/comp/image-atomic-automatic-bindings.comp
new file mode 100644
index 00000000..862cd212
--- /dev/null
+++ b/shaders-msl/comp/image-atomic-automatic-bindings.comp
@@ -0,0 +1,16 @@
+#version 450
+layout(local_size_x = 1) in;
+
+layout(set = 0, binding = 0, r32ui) uniform uimage2D uImage;
+layout(set = 0, binding = 1) uniform sampler2D uTexture;
+
+layout(set = 0, binding = 2) buffer SSBO
+{
+	vec4 outdata;
+};
+
+void main()
+{
+	uint ret = imageAtomicAdd(uImage, ivec2(gl_GlobalInvocationID.xy), 10u);
+	outdata = textureLod(uTexture, vec2(gl_GlobalInvocationID.xy), 0.0) + float(ret);
+}
diff --git a/spirv_msl.cpp b/spirv_msl.cpp
index b16e1e80..c0db4f2f 100644
--- a/spirv_msl.cpp
+++ b/spirv_msl.cpp
@@ -9627,9 +9627,14 @@ uint32_t CompilerMSL::get_metal_resource_index(SPIRVariable &var, SPIRType::Base
 	// If a matching binding has been specified, find and use it.
 	auto itr = resource_bindings.find({ execution.model, var_desc_set, var_binding });
 
-	auto resource_decoration = var_type.basetype == SPIRType::SampledImage && basetype == SPIRType::Sampler ?
-	                               SPIRVCrossDecorationResourceIndexSecondary :
-	                               SPIRVCrossDecorationResourceIndexPrimary;
+	// Atomic helper buffers for image atomics need to use secondary bindings as well.
+	bool use_secondary_binding = (var_type.basetype == SPIRType::SampledImage && basetype == SPIRType::Sampler) ||
+	                             basetype == SPIRType::AtomicCounter;
+
+	auto resource_decoration = use_secondary_binding ?
+	                           SPIRVCrossDecorationResourceIndexSecondary :
+	                           SPIRVCrossDecorationResourceIndexPrimary;
+
 	if (plane == 1)
 		resource_decoration = SPIRVCrossDecorationResourceIndexTertiary;
 	if (plane == 2)