SPIRV-Cross/reference/shaders-msl/comp/shared.comp

#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

struct SSBO
{
    float in_data[1];
};

struct SSBO2
{
    float out_data[1];
};

constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 1u, 1u);

kernel void main0(const device SSBO& _22 [[buffer(0)]], device SSBO2& _44 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
{
    threadgroup float sShared[4];
    uint ident = gl_GlobalInvocationID.x;
    float idata = _22.in_data[ident];
    sShared[gl_LocalInvocationIndex] = idata;
    threadgroup_barrier(mem_flags::mem_threadgroup);
    _44.out_data[ident] = sShared[(4u - gl_LocalInvocationIndex) - 1u];
}
Enhancements to MSL compute and entry point naming. Support Workgroup (threadgroup) variables. Mark if SPIRConstant is used as an array length, since it cannot be specialized. Resolve specialized array length constants. Support passing an array to MSL function. Support emitting GLSL array assignments in MSL via an array copy function. Support for memory and control barriers. Struct packing enhancements, including packing nested structs. Enhancements to replacing illegal MSL variable and function names. Add Compiler::get_entry_point_name_map() function to retrieve entry point renamings. Remove CompilerGLSL::clean_func_name() as obsolete. Fixes to types in bitcast MSL functions. Add Variant::get_id() member function. Add CompilerMSL::Options::msl_version option. Add numerous MSL compute tests. 2017-11-06 02:34:42 +00:00			`#include <metal_stdlib>`
			`#include <simd/simd.h>`

			`using namespace metal;`

			`struct SSBO`
			`{`
			`float in_data[1];`
			`};`

			`struct SSBO2`
			`{`
			`float out_data[1];`
			`};`

MSL: Declare gl_WorkGroupSize constant with [[maybe_unused]]. Avoids ugly warnings on nearly every compute shader. We could do analysis to detect whether we need to emit this constant, but it's a bit tedious to figure out if an OpConstantComponent is actually used by opcodes, so just make it simple. 2019-03-28 09:54:18 +00:00			`constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 1u, 1u);`
Rewrite how IDs are iterated over. This is a fairly fundamental change on how IDs are handled. It serves many purposes: - Improve performance. We only need to iterate over IDs which are relevant at any one time. - Makes sure we iterate through IDs in SPIR-V module declaration order rather than ID space. IDs don't have to be monotonically increasing, which was an assumption SPIRV-Cross used to have. It has apparently never been a problem until now. - Support LUTs of structs. We do this by interleaving declaration of constants and struct types in SPIR-V module order. To support this, the ParsedIR interface needed to change slightly. Before setting any ID with variant_set<T> we let ParsedIR know that an ID with a specific type has been added. The surface for change should be minimal. ParsedIR will maintain a per-type list of IDs which the cross-compiler will need to consider for later. Instead of looping over ir.ids[] (which can be extremely large), we loop over types now, using: ir.for_each_typed_id<SPIRVariable>([&](uint32_t id, SPIRVariable &var) { handle_variable(var); }); Now we make sure that we're never looking at irrelevant types. 2019-01-10 08:49:33 +00:00
Declare read-only SSBOs as const device in MSL. 2018-05-25 08:03:46 +00:00			`kernel void main0(const device SSBO& _22 [[buffer(0)]], device SSBO2& _44 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])`
Enhancements to MSL compute and entry point naming. Support Workgroup (threadgroup) variables. Mark if SPIRConstant is used as an array length, since it cannot be specialized. Resolve specialized array length constants. Support passing an array to MSL function. Support emitting GLSL array assignments in MSL via an array copy function. Support for memory and control barriers. Struct packing enhancements, including packing nested structs. Enhancements to replacing illegal MSL variable and function names. Add Compiler::get_entry_point_name_map() function to retrieve entry point renamings. Remove CompilerGLSL::clean_func_name() as obsolete. Fixes to types in bitcast MSL functions. Add Variant::get_id() member function. Add CompilerMSL::Options::msl_version option. Add numerous MSL compute tests. 2017-11-06 02:34:42 +00:00			`{`
Fix atomic_compare_exchange_weak_explicit. Need to emit a CAS loop. Fix shared memory declaration. Declare atomic ops with correct memory scope. 2018-05-15 14:03:20 +00:00			`threadgroup float sShared[4];`
Enhancements to MSL compute and entry point naming. Support Workgroup (threadgroup) variables. Mark if SPIRConstant is used as an array length, since it cannot be specialized. Resolve specialized array length constants. Support passing an array to MSL function. Support emitting GLSL array assignments in MSL via an array copy function. Support for memory and control barriers. Struct packing enhancements, including packing nested structs. Enhancements to replacing illegal MSL variable and function names. Add Compiler::get_entry_point_name_map() function to retrieve entry point renamings. Remove CompilerGLSL::clean_func_name() as obsolete. Fixes to types in bitcast MSL functions. Add Variant::get_id() member function. Add CompilerMSL::Options::msl_version option. Add numerous MSL compute tests. 2017-11-06 02:34:42 +00:00			`uint ident = gl_GlobalInvocationID.x;`
			`float idata = _22.in_data[ident];`
			`sShared[gl_LocalInvocationIndex] = idata;`
			`threadgroup_barrier(mem_flags::mem_threadgroup);`
			`_44.out_data[ident] = sShared[(4u - gl_LocalInvocationIndex) - 1u];`
			`}`