AuroraMiddleware/SPIRV-Cross
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity
fd738e3387
SPIRV-Cross/reference/shaders-msl/comp/culling.comp

37 lines
776 B
Plaintext
Raw Normal View History

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
#pragma clang diagnostic ignored "-Wunused-variable"
#include <metal_stdlib>
#include <simd/simd.h>
#include <metal_atomic>
using namespace metal;
struct SSBO
{
float in_data[1];
};
struct SSBO2
{
float out_data[1];
};
struct SSBO3
{
uint count;
};
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& _38 [[buffer(1)]], device SSBO3& _41 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
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];
if (idata > 12.0)
{
MSL: Handle coherent, volatile, and restrict. This maps them to their MSL equivalents. I've mapped `Coherent` to `volatile` since MSL doesn't have anything weaker than `volatile` but stronger than nothing. As part of this, I had to remove the implicit `volatile` added for atomic operation casts. If the buffer is already `coherent` or `volatile`, then we would add a second `volatile`, which would be redundant. I think this is OK even when the buffer *doesn't* have `coherent`: `T *` is implicitly convertible to `volatile T *`, but not vice-versa. It seems to compile OK at any rate. (Note that the non-`volatile` overloads of the atomic functions documented in the spec aren't present in the MSL 2.2 stdlib headers.) `restrict` is tricky, because in MSL, as in C++, it needs to go *after* the asterisk or ampersand for the pointer type it's modifying. Another issue is that, in the `Simple`, `GLSL450`, and `Vulkan` memory models, `Restrict` is the default (i.e. does not need to be specified); but MSL likely follows the `OpenCL` model where `Aliased` is the default. We probably need to implicitly set either `Restrict` or `Aliased` depending on the module's declared memory model.
2019-07-10 16:17:40 +00:00
uint _45 = atomic_fetch_add_explicit((device atomic_uint*)&_41.count, 1u, memory_order_relaxed);
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
_38.out_data[_45] = idata;
}
}
Reference in New Issue Copy Permalink