SPIRV-Cross/reference/shaders-msl/comp/coherent-image.comp

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

using namespace metal;

struct SSBO
{
    int4 value;
};

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

kernel void main0(volatile device SSBO& _10 [[buffer(0)]], texture2d<int> uImage [[texture(0)]])
{
    _10.value = uImage.read(uint2(int2(10)));
}
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`
			`{`
			`int4 value;`
			`};`

Update for pull request #1162 rev. 1 2019-09-24 22:13:04 +00:00			`constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);`

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			`kernel void main0(volatile device SSBO& _10 [[buffer(0)]], texture2d<int> uImage [[texture(0)]])`
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			`{`
			`_10.value = uImage.read(uint2(int2(10)));`
			`}`