Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Original file line number Diff line number Diff line change
@@ -0,0 +1,31 @@
#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

struct Record
{
uint a;
uint b;
uint c;
uint d;
};

struct WriteSSBO
{
Record result;
};

struct ReadSSBO
{
Record records[1];
};

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

kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device WriteSSBO& _10 [[buffer(0)]], device ReadSSBO& _16 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
constant uint& _10BufferSize = spvBufferSizeConstants[0];
constant uint& _16BufferSize = spvBufferSizeConstants[1];
_10.result = (uint(gl_GlobalInvocationID.x) < (_16BufferSize / 4u) ? _16.records[min(uint(gl_GlobalInvocationID.x), uint(max(int((_16BufferSize / 4u)) - 1, 0)))] : Record{});
}
Original file line number Diff line number Diff line change
@@ -0,0 +1,31 @@
#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

struct Record
{
uint a;
uint b;
uint c;
uint d;
};

struct WriteSSBO
{
Record result;
};

struct ReadSSBO
{
Record records[1];
};

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

kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device WriteSSBO& _10 [[buffer(0)]], device ReadSSBO& _16 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
constant uint& _10BufferSize = spvBufferSizeConstants[0];
constant uint& _16BufferSize = spvBufferSizeConstants[1];
_10.result = (uint(gl_GlobalInvocationID.x) < (_16BufferSize / 4u) ? _16.records[min(uint(gl_GlobalInvocationID.x), uint(max(int((_16BufferSize / 4u)) - 1, 0)))] : Record{});
}
Original file line number Diff line number Diff line change
@@ -0,0 +1,25 @@
#version 450
layout(local_size_x = 1) in;

struct Record
{
uint a;
uint b;
uint c;
uint d;
};

layout(std430, binding = 0) readonly buffer ReadSSBO
{
Record records[];
};

layout(std430, binding = 1) writeonly buffer WriteSSBO
{
Record result;
};

void main()
{
result = records[gl_GlobalInvocationID.x];
}
10 changes: 6 additions & 4 deletions spirv_msl.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -10336,12 +10336,14 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
if (expr)
{
auto &loaded_type = get<SPIRType>(result_type);
// Generate a proper zero expression using type constructor
// to_zero_initialized_expression returns {} which doesn't work in ternary
string zero_expr = join(type_to_glsl(loaded_type), "(0)");
// MSL structs do not have a scalar constructor. Use typed aggregate
// initialization so the expression remains valid in a ternary.
string zero_expr = loaded_type.basetype == SPIRType::Struct ? join(type_to_glsl(loaded_type), "{}") :
join(type_to_glsl(loaded_type), "(0)");

// Wrap the expression with a select: (in_bounds) ? loaded_value : zero
expr->expression = join("(", info.bounds_check_condition, " ? ", expr->expression, " : ", zero_expr, ")");
expr->expression =
join("(", info.bounds_check_condition, " ? ", expr->expression, " : ", zero_expr, ")");
}
// Remove from map to avoid double-processing
robust_access_chains.erase(it);
Expand Down