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

MSL: Also test image atomic for image2DArray.

Browse Source
This commit is contained in:
Hans-Kristian Arntzen 2024-02-26 15:43:29 +01:00
parent 7e0c9ab9ed
commit 855a5c369f
3 changed files with 93 additions and 89 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

89
reference/opt/shaders-msl/comp/atomic-image.msl31.comp
View File

@ -14,59 +14,60 @@ struct SSBO
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
kernel void main0(device SSBO& ssbo [[buffer(0)]], texture2d<uint, access::read_write> uImage [[texture(0)]], texture2d<int, access::read_write> iImage [[texture(1)]])
kernel void main0(device SSBO& ssbo [[buffer(0)]], texture2d<uint, access::read_write> uImage [[texture(0)]], texture2d_array<uint, access::read_write> uImageArray [[texture(1)]], texture2d<int, access::read_write> iImage [[texture(2)]])
{
uint _19 = uImage.atomic_fetch_add(uint2(int2(1, 5)), 1u).x;
uint _27 = uImage.atomic_fetch_add(uint2(int2(1, 5)), 1u).x;
iImage.write(int4(int(_27)), uint2(int2(1, 6)));
uint _32 = uImage.atomic_fetch_or(uint2(int2(1, 5)), 1u).x;
uint _34 = uImage.atomic_fetch_xor(uint2(int2(1, 5)), 1u).x;
uint _36 = uImage.atomic_fetch_and(uint2(int2(1, 5)), 1u).x;
uint _38 = uImage.atomic_fetch_min(uint2(int2(1, 5)), 1u).x;
uint _40 = uImage.atomic_fetch_max(uint2(int2(1, 5)), 1u).x;
uint _44;
uint4 _102;
uint _27 = uImageArray.atomic_fetch_add(uint3(int3(1, 5, 8)).xy, uint3(int3(1, 5, 8)).z, 1u).x;
uint _35 = uImage.atomic_fetch_add(uint2(int2(1, 5)), 1u).x;
iImage.write(int4(int(_35)), uint2(int2(1, 6)));
uint _40 = uImage.atomic_fetch_or(uint2(int2(1, 5)), 1u).x;
uint _42 = uImage.atomic_fetch_xor(uint2(int2(1, 5)), 1u).x;
uint _44 = uImage.atomic_fetch_and(uint2(int2(1, 5)), 1u).x;
uint _46 = uImage.atomic_fetch_min(uint2(int2(1, 5)), 1u).x;
uint _48 = uImage.atomic_fetch_max(uint2(int2(1, 5)), 1u).x;
uint _52;
uint4 _110;
do
{
_102.x = 10u;
} while (!uImage.atomic_compare_exchange_weak(uint2(int2(1, 5)), &_102, 2u) && _102.x == 10u);
_44 = _102.x;
int _47 = iImage.atomic_fetch_add(uint2(int2(1, 6)), 1).x;
int _49 = iImage.atomic_fetch_or(uint2(int2(1, 6)), 1).x;
int _51 = iImage.atomic_fetch_xor(uint2(int2(1, 6)), 1).x;
int _53 = iImage.atomic_fetch_and(uint2(int2(1, 6)), 1).x;
int _55 = iImage.atomic_fetch_min(uint2(int2(1, 6)), 1).x;
int _57 = iImage.atomic_fetch_max(uint2(int2(1, 6)), 1).x;
int _61;
int4 _104;
_110.x = 10u;
} while (!uImage.atomic_compare_exchange_weak(uint2(int2(1, 5)), &_110, 2u) && _110.x == 10u);
_52 = _110.x;
int _55 = iImage.atomic_fetch_add(uint2(int2(1, 6)), 1).x;
int _57 = iImage.atomic_fetch_or(uint2(int2(1, 6)), 1).x;
int _59 = iImage.atomic_fetch_xor(uint2(int2(1, 6)), 1).x;
int _61 = iImage.atomic_fetch_and(uint2(int2(1, 6)), 1).x;
int _63 = iImage.atomic_fetch_min(uint2(int2(1, 6)), 1).x;
int _65 = iImage.atomic_fetch_max(uint2(int2(1, 6)), 1).x;
int _69;
int4 _112;
do
{
_104.x = 10;
} while (!iImage.atomic_compare_exchange_weak(uint2(int2(1, 5)), &_104, 2) && _104.x == 10);
_61 = _104.x;
uint _68 = atomic_fetch_add_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _70 = atomic_fetch_or_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _72 = atomic_fetch_xor_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _74 = atomic_fetch_and_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _76 = atomic_fetch_min_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _78 = atomic_fetch_max_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _80 = atomic_exchange_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _82;
_112.x = 10;
} while (!iImage.atomic_compare_exchange_weak(uint2(int2(1, 5)), &_112, 2) && _112.x == 10);
_69 = _112.x;
uint _76 = atomic_fetch_add_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _78 = atomic_fetch_or_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _80 = atomic_fetch_xor_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _82 = atomic_fetch_and_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _84 = atomic_fetch_min_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _86 = atomic_fetch_max_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _88 = atomic_exchange_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _90;
do
{
_82 = 10u;
} while (!atomic_compare_exchange_weak_explicit((device atomic_uint*)&ssbo.u32, &_82, 2u, memory_order_relaxed, memory_order_relaxed) && _82 == 10u);
int _85 = atomic_fetch_add_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _87 = atomic_fetch_or_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _89 = atomic_fetch_xor_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _91 = atomic_fetch_and_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _93 = atomic_fetch_min_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _95 = atomic_fetch_max_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _97 = atomic_exchange_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _99;
_90 = 10u;
} while (!atomic_compare_exchange_weak_explicit((device atomic_uint*)&ssbo.u32, &_90, 2u, memory_order_relaxed, memory_order_relaxed) && _90 == 10u);
int _93 = atomic_fetch_add_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _95 = atomic_fetch_or_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _97 = atomic_fetch_xor_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _99 = atomic_fetch_and_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _101 = atomic_fetch_min_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _103 = atomic_fetch_max_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _105 = atomic_exchange_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _107;
do
{
_99 = 10;
} while (!atomic_compare_exchange_weak_explicit((device atomic_int*)&ssbo.i32, &_99, 2, memory_order_relaxed, memory_order_relaxed) && _99 == 10);
_107 = 10;
} while (!atomic_compare_exchange_weak_explicit((device atomic_int*)&ssbo.i32, &_107, 2, memory_order_relaxed, memory_order_relaxed) && _107 == 10);
}

89
reference/shaders-msl/comp/atomic-image.msl31.comp
View File

@ -14,59 +14,60 @@ struct SSBO
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
kernel void main0(device SSBO& ssbo [[buffer(0)]], texture2d<uint, access::read_write> uImage [[texture(0)]], texture2d<int, access::read_write> iImage [[texture(1)]])
kernel void main0(device SSBO& ssbo [[buffer(0)]], texture2d<uint, access::read_write> uImage [[texture(0)]], texture2d_array<uint, access::read_write> uImageArray [[texture(1)]], texture2d<int, access::read_write> iImage [[texture(2)]])
{
uint _19 = uImage.atomic_fetch_add(uint2(int2(1, 5)), 1u).x;
uint _27 = uImage.atomic_fetch_add(uint2(int2(1, 5)), 1u).x;
iImage.write(int4(int(_27)), uint2(int2(1, 6)));
uint _32 = uImage.atomic_fetch_or(uint2(int2(1, 5)), 1u).x;
uint _34 = uImage.atomic_fetch_xor(uint2(int2(1, 5)), 1u).x;
uint _36 = uImage.atomic_fetch_and(uint2(int2(1, 5)), 1u).x;
uint _38 = uImage.atomic_fetch_min(uint2(int2(1, 5)), 1u).x;
uint _40 = uImage.atomic_fetch_max(uint2(int2(1, 5)), 1u).x;
uint _44;
uint4 _102;
uint _27 = uImageArray.atomic_fetch_add(uint3(int3(1, 5, 8)).xy, uint3(int3(1, 5, 8)).z, 1u).x;
uint _35 = uImage.atomic_fetch_add(uint2(int2(1, 5)), 1u).x;
iImage.write(int4(int(_35)), uint2(int2(1, 6)));
uint _40 = uImage.atomic_fetch_or(uint2(int2(1, 5)), 1u).x;
uint _42 = uImage.atomic_fetch_xor(uint2(int2(1, 5)), 1u).x;
uint _44 = uImage.atomic_fetch_and(uint2(int2(1, 5)), 1u).x;
uint _46 = uImage.atomic_fetch_min(uint2(int2(1, 5)), 1u).x;
uint _48 = uImage.atomic_fetch_max(uint2(int2(1, 5)), 1u).x;
uint _52;
uint4 _110;
do
{
_102.x = 10u;
} while (!uImage.atomic_compare_exchange_weak(uint2(int2(1, 5)), &_102, 2u) && _102.x == 10u);
_44 = _102.x;
int _47 = iImage.atomic_fetch_add(uint2(int2(1, 6)), 1).x;
int _49 = iImage.atomic_fetch_or(uint2(int2(1, 6)), 1).x;
int _51 = iImage.atomic_fetch_xor(uint2(int2(1, 6)), 1).x;
int _53 = iImage.atomic_fetch_and(uint2(int2(1, 6)), 1).x;
int _55 = iImage.atomic_fetch_min(uint2(int2(1, 6)), 1).x;
int _57 = iImage.atomic_fetch_max(uint2(int2(1, 6)), 1).x;
int _61;
int4 _104;
_110.x = 10u;
} while (!uImage.atomic_compare_exchange_weak(uint2(int2(1, 5)), &_110, 2u) && _110.x == 10u);
_52 = _110.x;
int _55 = iImage.atomic_fetch_add(uint2(int2(1, 6)), 1).x;
int _57 = iImage.atomic_fetch_or(uint2(int2(1, 6)), 1).x;
int _59 = iImage.atomic_fetch_xor(uint2(int2(1, 6)), 1).x;
int _61 = iImage.atomic_fetch_and(uint2(int2(1, 6)), 1).x;
int _63 = iImage.atomic_fetch_min(uint2(int2(1, 6)), 1).x;
int _65 = iImage.atomic_fetch_max(uint2(int2(1, 6)), 1).x;
int _69;
int4 _112;
do
{
_104.x = 10;
} while (!iImage.atomic_compare_exchange_weak(uint2(int2(1, 5)), &_104, 2) && _104.x == 10);
_61 = _104.x;
uint _68 = atomic_fetch_add_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _70 = atomic_fetch_or_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _72 = atomic_fetch_xor_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _74 = atomic_fetch_and_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _76 = atomic_fetch_min_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _78 = atomic_fetch_max_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _80 = atomic_exchange_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _82;
_112.x = 10;
} while (!iImage.atomic_compare_exchange_weak(uint2(int2(1, 5)), &_112, 2) && _112.x == 10);
_69 = _112.x;
uint _76 = atomic_fetch_add_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _78 = atomic_fetch_or_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _80 = atomic_fetch_xor_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _82 = atomic_fetch_and_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _84 = atomic_fetch_min_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _86 = atomic_fetch_max_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _88 = atomic_exchange_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
uint _90;
do
{
_82 = 10u;
} while (!atomic_compare_exchange_weak_explicit((device atomic_uint*)&ssbo.u32, &_82, 2u, memory_order_relaxed, memory_order_relaxed) && _82 == 10u);
int _85 = atomic_fetch_add_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _87 = atomic_fetch_or_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _89 = atomic_fetch_xor_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _91 = atomic_fetch_and_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _93 = atomic_fetch_min_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _95 = atomic_fetch_max_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _97 = atomic_exchange_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _99;
_90 = 10u;
} while (!atomic_compare_exchange_weak_explicit((device atomic_uint*)&ssbo.u32, &_90, 2u, memory_order_relaxed, memory_order_relaxed) && _90 == 10u);
int _93 = atomic_fetch_add_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _95 = atomic_fetch_or_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _97 = atomic_fetch_xor_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _99 = atomic_fetch_and_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _101 = atomic_fetch_min_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _103 = atomic_fetch_max_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _105 = atomic_exchange_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _107;
do
{
_99 = 10;
} while (!atomic_compare_exchange_weak_explicit((device atomic_int*)&ssbo.i32, &_99, 2, memory_order_relaxed, memory_order_relaxed) && _99 == 10);
_107 = 10;
} while (!atomic_compare_exchange_weak_explicit((device atomic_int*)&ssbo.i32, &_107, 2, memory_order_relaxed, memory_order_relaxed) && _107 == 10);
}

4
shaders-msl/comp/atomic-image.msl31.comp
View File

@ -4,7 +4,8 @@ layout(local_size_x = 1) in;
layout(r32ui, binding = 0) uniform highp uimage2D uImage;
layout(r32i, binding = 1) uniform highp iimage2D iImage;
layout(binding = 2, std430) buffer SSBO
layout(r32ui, binding = 2) uniform highp uimage2DArray uImageArray;
layout(binding = 3, std430) buffer SSBO
{
uint u32;
int i32;
@ -13,6 +14,7 @@ layout(binding = 2, std430) buffer SSBO
void main()
{
imageAtomicAdd(uImage, ivec2(1, 5), 1u);
imageAtomicAdd(uImageArray, ivec3(1, 5, 8), 1u);
// Test that we do not invalidate OpImage variables which are loaded from UniformConstant
// address space.