GLSL/MSL: Support VK_KHR_zero_initialize_workgroup_memory

reference/opt/shaders-msl/comp/shared-zero-init-simple.comp

@@ -0,0 +1,31 @@
+#pragma clang diagnostic ignored "-Wsometimes-uninitialized"
+#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& _32 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
+{
+    threadgroup float sShared;
+    {
+        if (gl_LocalInvocationIndex == 0)
+        {
+            sShared = 0.0;
+        }
+        threadgroup_barrier(mem_flags::mem_threadgroup);
+    }
+    _32.out_data[gl_GlobalInvocationID.x] = sShared + _22.in_data[gl_GlobalInvocationID.x];
+}
+

reference/opt/shaders-msl/comp/shared-zero-init.comp

@@ -0,0 +1,90 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+    T elements[Num ? Num : 1];
+
+    thread T& operator [] (size_t pos) thread
+    {
+        return elements[pos];
+    }
+    constexpr const thread T& operator [] (size_t pos) const thread
+    {
+        return elements[pos];
+    }
+
+    device T& operator [] (size_t pos) device
+    {
+        return elements[pos];
+    }
+    constexpr const device T& operator [] (size_t pos) const device
+    {
+        return elements[pos];
+    }
+
+    constexpr const constant T& operator [] (size_t pos) const constant
+    {
+        return elements[pos];
+    }
+
+    threadgroup T& operator [] (size_t pos) threadgroup
+    {
+        return elements[pos];
+    }
+    constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+    {
+        return elements[pos];
+    }
+};
+
+struct SSBO
+{
+    float in_data[1];
+};
+
+struct SSBO2
+{
+    float out_data[1];
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 1u, 1u);
+
+constant spvUnsafeArray<float, 4> _31 = spvUnsafeArray<float, 4>({ 0.0, 0.0, 0.0, 0.0 });
+
+kernel void main0(const device SSBO& _22 [[buffer(0)]], device SSBO2& _48 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
+{
+    threadgroup spvUnsafeArray<float, 4> sShared;
+    {
+        threadgroup uint *sShared_ptr = (threadgroup uint *)&sShared;
+        uint sShared_sz = sizeof(sShared);
+        uint sShared_pos = gl_LocalInvocationIndex;
+        uint sShared_stride = gl_WorkGroupSize.x * gl_WorkGroupSize.y * gl_WorkGroupSize.z;
+        while (sizeof(uint) * sShared_pos < sShared_sz)
+        {
+            sShared_ptr[sShared_pos] = 0u;
+            sShared_pos += sShared_stride;
+        }
+        if (gl_LocalInvocationIndex == 0)
+        {
+            sShared_pos = (sShared_sz / sizeof(uint)) * sizeof(uint);
+            threadgroup uchar *sShared_ptr2 = (threadgroup uchar *)&sShared;
+            while (sShared_pos < sShared_sz)
+            {
+                sShared_ptr2[sShared_pos] = '\0';
+                sShared_pos++;
+            }
+        }
+        threadgroup_barrier(mem_flags::mem_threadgroup);
+    }
+    sShared[gl_LocalInvocationIndex] += _22.in_data[gl_GlobalInvocationID.x];
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    _48.out_data[gl_GlobalInvocationID.x] = sShared[3u - gl_LocalInvocationIndex];
+}
+

reference/opt/shaders/comp/spec-const-arraydim-init.comp

@@ -0,0 +1,27 @@
+#version 450
+#extension GL_EXT_null_initializer : require
+layout(local_size_x = 2, local_size_y = 1, local_size_z = 1) in;
+
+struct Data
+{
+    float a;
+    float b;
+};
+
+#ifndef SPIRV_CROSS_CONSTANT_ID_0
+#define SPIRV_CROSS_CONSTANT_ID_0 2
+#endif
+const int arraySize = SPIRV_CROSS_CONSTANT_ID_0;
+const Data _25[arraySize] = { };
+
+layout(binding = 0, std430) buffer SSBO
+{
+    Data outdata[];
+} _11;
+
+void main()
+{
+    _11.outdata[gl_WorkGroupID.x].a = _25[gl_WorkGroupID.x].a;
+    _11.outdata[gl_WorkGroupID.x].b = _25[gl_WorkGroupID.x].b;
+}
+

reference/shaders-msl/comp/shared-zero-init-simple.comp

@@ -0,0 +1,33 @@
+#pragma clang diagnostic ignored "-Wsometimes-uninitialized"
+#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& _32 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
+{
+    threadgroup float sShared;
+    {
+        if (gl_LocalInvocationIndex == 0)
+        {
+            sShared = 0.0;
+        }
+        threadgroup_barrier(mem_flags::mem_threadgroup);
+    }
+    uint ident = gl_GlobalInvocationID.x;
+    float idata = _22.in_data[ident];
+    _32.out_data[ident] = sShared + idata;
+}
+

reference/shaders-msl/comp/shared-zero-init.comp

@@ -0,0 +1,92 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+    T elements[Num ? Num : 1];
+
+    thread T& operator [] (size_t pos) thread
+    {
+        return elements[pos];
+    }
+    constexpr const thread T& operator [] (size_t pos) const thread
+    {
+        return elements[pos];
+    }
+
+    device T& operator [] (size_t pos) device
+    {
+        return elements[pos];
+    }
+    constexpr const device T& operator [] (size_t pos) const device
+    {
+        return elements[pos];
+    }
+
+    constexpr const constant T& operator [] (size_t pos) const constant
+    {
+        return elements[pos];
+    }
+
+    threadgroup T& operator [] (size_t pos) threadgroup
+    {
+        return elements[pos];
+    }
+    constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+    {
+        return elements[pos];
+    }
+};
+
+struct SSBO
+{
+    float in_data[1];
+};
+
+struct SSBO2
+{
+    float out_data[1];
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 1u, 1u);
+
+constant spvUnsafeArray<float, 4> _31 = spvUnsafeArray<float, 4>({ 0.0, 0.0, 0.0, 0.0 });
+
+kernel void main0(const device SSBO& _22 [[buffer(0)]], device SSBO2& _48 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
+{
+    threadgroup spvUnsafeArray<float, 4> sShared;
+    {
+        threadgroup uint *sShared_ptr = (threadgroup uint *)&sShared;
+        uint sShared_sz = sizeof(sShared);
+        uint sShared_pos = gl_LocalInvocationIndex;
+        uint sShared_stride = gl_WorkGroupSize.x * gl_WorkGroupSize.y * gl_WorkGroupSize.z;
+        while (sizeof(uint) * sShared_pos < sShared_sz)
+        {
+            sShared_ptr[sShared_pos] = 0u;
+            sShared_pos += sShared_stride;
+        }
+        if (gl_LocalInvocationIndex == 0)
+        {
+            sShared_pos = (sShared_sz / sizeof(uint)) * sizeof(uint);
+            threadgroup uchar *sShared_ptr2 = (threadgroup uchar *)&sShared;
+            while (sShared_pos < sShared_sz)
+            {
+                sShared_ptr2[sShared_pos] = '\0';
+                sShared_pos++;
+            }
+        }
+        threadgroup_barrier(mem_flags::mem_threadgroup);
+    }
+    uint ident = gl_GlobalInvocationID.x;
+    float idata = _22.in_data[ident];
+    sShared[gl_LocalInvocationIndex] += idata;
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    _48.out_data[ident] = sShared[(4u - gl_LocalInvocationIndex) - 1u];
+}
+

reference/shaders/comp/spec-const-arraydim-init.comp

@@ -0,0 +1,27 @@
+#version 450
+#extension GL_EXT_null_initializer : require
+layout(local_size_x = 2, local_size_y = 1, local_size_z = 1) in;
+
+struct Data
+{
+    float a;
+    float b;
+};
+
+#ifndef SPIRV_CROSS_CONSTANT_ID_0
+#define SPIRV_CROSS_CONSTANT_ID_0 2
+#endif
+const int arraySize = SPIRV_CROSS_CONSTANT_ID_0;
+const Data _25[arraySize] = { };
+
+layout(binding = 0, std430) buffer SSBO
+{
+    Data outdata[];
+} _11;
+
+void main()
+{
+    _11.outdata[gl_WorkGroupID.x].a = _25[gl_WorkGroupID.x].a;
+    _11.outdata[gl_WorkGroupID.x].b = _25[gl_WorkGroupID.x].b;
+}
+

shaders-msl/comp/shared-zero-init-simple.comp

@@ -0,0 +1,24 @@
+#version 450
+#extension GL_EXT_null_initializer : enable
+layout(local_size_x = 4) in;
+
+shared float sShared = {};
+
+layout(std430, binding = 0) readonly buffer SSBO
+{
+    float in_data[];
+};
+
+layout(std430, binding = 1) writeonly buffer SSBO2
+{
+    float out_data[];
+};
+
+void main()
+{
+    uint ident = gl_GlobalInvocationID.x;
+    float idata = in_data[ident];
+
+    out_data[ident] = sShared + idata;
+}
+

shaders-msl/comp/shared-zero-init.comp

@@ -0,0 +1,28 @@
+#version 450
+#extension GL_EXT_null_initializer : enable
+layout(local_size_x = 4) in;
+
+shared float sShared[gl_WorkGroupSize.x] = {};
+
+layout(std430, binding = 0) readonly buffer SSBO
+{
+    float in_data[];
+};
+
+layout(std430, binding = 1) writeonly buffer SSBO2
+{
+    float out_data[];
+};
+
+void main()
+{
+    uint ident = gl_GlobalInvocationID.x;
+    float idata = in_data[ident];
+
+    sShared[gl_LocalInvocationIndex] += idata;
+    memoryBarrierShared();
+    barrier();
+
+    out_data[ident] = sShared[gl_WorkGroupSize.x - gl_LocalInvocationIndex - 1u];
+}
+

shaders/comp/spec-const-arraydim-init.comp

@@ -0,0 +1,22 @@
+#version 450
+#extension GL_EXT_null_initializer : require
+
+layout(constant_id = 0) const int arraySize = 2;
+layout(local_size_x = 2) in;
+
+struct Data
+{
+	float a;
+	float b;
+};
+
+layout(std430, binding = 0) buffer SSBO
+{
+	Data outdata[];
+};
+
+void main()
+{
+	Data d[arraySize] = {};
+	outdata[gl_WorkGroupID.x] = d[gl_WorkGroupID.x];
+}

spirv_common.hpp

@@ -1410,6 +1410,10 @@ struct SPIRConstant : IVariant
 	// If true, this is a LUT, and should always be declared in the outer scope.
 	bool is_used_as_lut = false;
 
+	// If this is a null constant of array type with specialized length.
+	// May require special handling in initializer
+	bool is_null_array_specialized_length = false;
+
 	// For composites which are constant arrays, etc.
 	SmallVector<ConstantID> subconstants;