8f6df770ce
Add bool members is_read and is_written to SPIRType::Image. Output correct texture read/write access by marking whether textures are read from and written to by the shader. Override bitcast_glsl_op() to use Metal as_type<type> functions. Add implementations of SPIR-V functions inverse(), degrees() & radians(). Map inverseSqrt() to rsqrt(). Map roundEven() to rint(). GLSL functions imageSize() and textureSize() map to equivalent expression using MSL get_width() & get_height() functions. Map several SPIR-V integer bitfield functions to MSL equivalents. Map SPIR-V atomic functions to MSL equivalents. Map texture packing and unpacking functions to MSL equivalents. Refactor existing, and add new, image query functions. Reorganize header lines into includes and pragmas. Simplify type_to_glsl() logic. Add MSL test case vert/functions.vert for added function implementations. Add MSL test case comp/atomic.comp for added function implementations. test_shaders.py use macOS compilation for MSL shader compilation validations.
37 lines
2.1 KiB
Plaintext
37 lines
2.1 KiB
Plaintext
#pragma clang diagnostic ignored "-Wunused-variable"
|
|
|
|
#include <metal_stdlib>
|
|
#include <simd/simd.h>
|
|
#include <metal_atomic>
|
|
|
|
using namespace metal;
|
|
|
|
struct SSBO
|
|
{
|
|
uint u32;
|
|
int i32;
|
|
};
|
|
|
|
kernel void main0(device SSBO& ssbo [[buffer(0)]])
|
|
{
|
|
uint _16 = atomic_fetch_add_explicit((volatile device atomic_uint*)&(ssbo.u32), 1u, memory_order_relaxed);
|
|
uint _18 = atomic_fetch_or_explicit((volatile device atomic_uint*)&(ssbo.u32), 1u, memory_order_relaxed);
|
|
uint _20 = atomic_fetch_xor_explicit((volatile device atomic_uint*)&(ssbo.u32), 1u, memory_order_relaxed);
|
|
uint _22 = atomic_fetch_and_explicit((volatile device atomic_uint*)&(ssbo.u32), 1u, memory_order_relaxed);
|
|
uint _24 = atomic_fetch_min_explicit((volatile device atomic_uint*)&(ssbo.u32), 1u, memory_order_relaxed);
|
|
uint _26 = atomic_fetch_max_explicit((volatile device atomic_uint*)&(ssbo.u32), 1u, memory_order_relaxed);
|
|
uint _28 = atomic_exchange_explicit((volatile device atomic_uint*)&(ssbo.u32), 1u, memory_order_relaxed);
|
|
uint _30 = 10u;
|
|
uint _32 = atomic_compare_exchange_weak_explicit((volatile device atomic_uint*)&(ssbo.u32), &(_30), 2u, memory_order_relaxed, memory_order_relaxed);
|
|
int _36 = atomic_fetch_add_explicit((volatile device atomic_int*)&(ssbo.i32), 1, memory_order_relaxed);
|
|
int _38 = atomic_fetch_or_explicit((volatile device atomic_int*)&(ssbo.i32), 1, memory_order_relaxed);
|
|
int _40 = atomic_fetch_xor_explicit((volatile device atomic_int*)&(ssbo.i32), 1, memory_order_relaxed);
|
|
int _42 = atomic_fetch_and_explicit((volatile device atomic_int*)&(ssbo.i32), 1, memory_order_relaxed);
|
|
int _44 = atomic_fetch_min_explicit((volatile device atomic_int*)&(ssbo.i32), 1, memory_order_relaxed);
|
|
int _46 = atomic_fetch_max_explicit((volatile device atomic_int*)&(ssbo.i32), 1, memory_order_relaxed);
|
|
int _48 = atomic_exchange_explicit((volatile device atomic_int*)&(ssbo.i32), 1, memory_order_relaxed);
|
|
int _50 = 10;
|
|
int _52 = atomic_compare_exchange_weak_explicit((volatile device atomic_int*)&(ssbo.i32), &(_50), 2, memory_order_relaxed, memory_order_relaxed);
|
|
}
|
|
|