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

msl: Add support for gl_LocalInvocationID and gl_LocalInvocationIndex

Browse Source
This commit is contained in:
Corentin Wallez 2017-01-26 20:22:31 -05:00
parent 3542168f51
commit d8278a8e3b
5 changed files with 78 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

24
reference/shaders-msl/comp/local-invocation-id.comp Normal file
View File

@ -0,0 +1,24 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct myBlock
{
int a;
float b[1];
};
// Support GLSL mod(), which is slightly different than Metal fmod()
template<typename Tx, typename Ty>
Tx mod(Tx x, Ty y)
{
return x - y * floor(x / y);
}
kernel void main0(device myBlock& myStorage [[buffer(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
{
myStorage.a = (myStorage.a + 1) % 256;
myStorage.b[gl_LocalInvocationID.x] = mod(myStorage.b[gl_LocalInvocationID.x] + 0.0199999995529651641845703125, 1.0);
}

24
reference/shaders-msl/comp/local-invocation-index.comp Normal file
View File

@ -0,0 +1,24 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct myBlock
{
int a;
float b[1];
};
// Support GLSL mod(), which is slightly different than Metal fmod()
template<typename Tx, typename Ty>
Tx mod(Tx x, Ty y)
{
return x - y * floor(x / y);
}
kernel void main0(device myBlock& myStorage [[buffer(0)]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
{
myStorage.a = (myStorage.a + 1) % 256;
myStorage.b[gl_LocalInvocationIndex] = mod(myStorage.b[gl_LocalInvocationIndex] + 0.0199999995529651641845703125, 1.0);
}

9
shaders-msl/comp/local-invocation-id.comp Normal file
View File

@ -0,0 +1,9 @@
#version 450
layout(set = 0, binding = 0) buffer myBlock {
int a;
float b[1];
} myStorage;
void main() {
myStorage.a = (myStorage.a + 1) % 256;
myStorage.b[gl_LocalInvocationID.x] = mod((myStorage.b[gl_LocalInvocationID.x] + 0.02), 1.0);
}

9
shaders-msl/comp/local-invocation-index.comp Normal file
View File

@ -0,0 +1,9 @@
#version 450
layout(set = 0, binding = 0) buffer myBlock {
int a;
float b[1];
} myStorage;
void main() {
myStorage.a = (myStorage.a + 1) % 256;
myStorage.b[gl_LocalInvocationIndex.x] = mod((myStorage.b[gl_LocalInvocationIndex.x] + 0.02), 1.0);
}

12
spirv_msl.cpp
View File

@ -1108,6 +1108,8 @@ string CompilerMSL::member_attribute_qualifier(const SPIRType &type, uint32_t in
switch (builtin)
{
case BuiltInGlobalInvocationId:
case BuiltInLocalInvocationId:
case BuiltInLocalInvocationIndex:
return string(" [[") + builtin_qualifier(builtin) + "]]";
default:
@ -1640,6 +1642,12 @@ string CompilerMSL::builtin_qualifier(BuiltIn builtin)
case BuiltInGlobalInvocationId:
return "thread_position_in_grid";
case BuiltInLocalInvocationId:
return "thread_position_in_threadgroup";
case BuiltInLocalInvocationIndex:
return "thread_index_in_threadgroup";
default:
return "unsupported-built-in";
}
@ -1683,6 +1691,10 @@ string CompilerMSL::builtin_type_decl(BuiltIn builtin)
// Compute function in
case BuiltInGlobalInvocationId:
return "uint3";
case BuiltInLocalInvocationId:
return "uint3";
case BuiltInLocalInvocationIndex:
return "uint";
default:
return "unsupported-built-in-type";