|
// Methodology: start with pctl_kernel.hlsl from https://github.com/bzm3r/piet-dx12 |
|
// Expand includes by hand. |
|
// Use http://shader-playground.timjones.io/ |
|
// Use DXC compiler, cs_6_0 profile, set entry point, SPIRV out |
|
// Chain spirv-cross, Metal out |
|
|
|
#include <metal_stdlib> |
|
#include <simd/simd.h> |
|
|
|
using namespace metal; |
|
|
|
struct type_Constants |
|
{ |
|
uint num_objects; |
|
uint object_size; |
|
uint tile_size; |
|
uint num_tiles_x; |
|
uint num_tiles_y; |
|
}; |
|
|
|
struct type_ByteAddressBuffer |
|
{ |
|
uint _m0[1]; |
|
}; |
|
|
|
struct type_RWByteAddressBuffer |
|
{ |
|
uint _m0[1]; |
|
}; |
|
|
|
kernel void build_per_tile_command_list(constant type_Constants& Constants [[buffer(0)]], const device type_ByteAddressBuffer& object_data_buffer [[buffer(1)]], device type_RWByteAddressBuffer& per_tile_command_list [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) |
|
{ |
|
uint _52 = (4u + (Constants.num_objects * Constants.object_size)) * ((Constants.num_tiles_x * gl_GlobalInvocationID.y) + gl_GlobalInvocationID.x); |
|
uint _56 = Constants.tile_size * gl_GlobalInvocationID.x; |
|
uint _57 = Constants.tile_size * gl_GlobalInvocationID.y; |
|
uint _61; |
|
_61 = 0u; |
|
uint _62; |
|
for (uint _64 = 0u; _64 < Constants.num_objects; _61 = _62, _64++) |
|
{ |
|
uint _69 = _64 * 24u; |
|
uint _72 = (_69 + 12u) >> 2u; |
|
uint _74 = object_data_buffer._m0[_72]; |
|
uint _75 = (_69 + 16u) >> 2u; |
|
uint _77 = object_data_buffer._m0[_75]; |
|
bool _87 = (((_56 + Constants.tile_size) <= ((4294901760u & _74) >> 16u)) || (_56 >= (65535u & _74))) ? false : true; |
|
if ((_87 && (((_57 + Constants.tile_size) <= ((4294901760u & _77) >> 16u)) || (_57 >= (65535u & _77)))) ? false : _87) |
|
{ |
|
uint _101 = object_data_buffer._m0[(_69 + 4u) >> 2u]; |
|
uint _104 = object_data_buffer._m0[(_69 + 8u) >> 2u]; |
|
uint _108 = object_data_buffer._m0[(_69 + 20u) >> 2u]; |
|
uint _110 = (_52 + 4u) + (_61 * Constants.object_size); |
|
per_tile_command_list._m0[_110 >> 2u] = object_data_buffer._m0[_69 >> 2u]; |
|
uint _114 = (_110 + 4u) >> 2u; |
|
per_tile_command_list._m0[_114] = _101; |
|
per_tile_command_list._m0[_114 + 1u] = _104; |
|
uint _119 = (_110 + 12u) >> 2u; |
|
per_tile_command_list._m0[_119] = _74; |
|
per_tile_command_list._m0[_119 + 1u] = _77; |
|
per_tile_command_list._m0[(_110 + 20u) >> 2u] = _108; |
|
_62 = _61 + 1u; |
|
continue; |
|
} |
|
else |
|
{ |
|
_62 = _61; |
|
continue; |
|
} |
|
continue; |
|
} |
|
per_tile_command_list._m0[_52 >> 2u] = _61; |
|
} |
|
|
