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

Merge pull request #1221 from KhronosGroup/fix-1219

Browse Source 
MSL: Fix automatic binding allocation for image atomic buffers.
This commit is contained in:
Hans-Kristian Arntzen 2019-11-28 13:21:03 +01:00 committed by GitHub
commit e46c8937ee
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
8 changed files with 79 additions and 7 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

25
reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.comp Normal file
View File

@ -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));
}

2
reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.frag
View File

@ -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);

2
reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.graphics-robust-access.frag
View File

@ -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];

26
reference/shaders-msl/comp/image-atomic-automatic-bindings.comp Normal file
View File

@ -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));
}

2
reference/shaders-ue4/asm/frag/texture-atomics.asm.frag
View File

@ -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);

2
reference/shaders-ue4/asm/frag/texture-atomics.asm.graphics-robust-access.frag
View File

@ -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);

16
shaders-msl/comp/image-atomic-automatic-bindings.comp Normal file
View File

@ -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);
}

11
spirv_msl.cpp
View File

@ -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)