rafx-framework 0.0.16

2��Wy�
Bstatic const uint3 gl_WorkGroupSize = uint3(64u, 1u, 1u);

cbuffer ClearBufferConfig : register(b0, space0)
{
    uint config_buffer_bytes_div_by_four : packoffset(c0);
    uint config_fill_value : packoffset(c0.y);
    uint config_num_workgroups_x : packoffset(c0.z);
};

RWByteAddressBuffer data : register(u1, space0);

static uint3 gl_GlobalInvocationID;
struct SPIRV_Cross_Input
{
    uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
};

void comp_main()
{
    uint index = (gl_GlobalInvocationID.y * (config_num_workgroups_x * 64u)) + gl_GlobalInvocationID.x;
    if (index < config_buffer_bytes_div_by_four)
    {
        data.Store(index * 4 + 0, config_fill_value);
    }
}

[numthreads(64, 1, 1)]
void main(SPIRV_Cross_Input stage_input)
{
    gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID;
    comp_main();
}

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

using namespace metal;

struct ClearBufferConfig
{
    uint buffer_bytes_div_by_four;
    uint fill_value;
    uint num_workgroups_x;
};

struct Buffer
{
    uint data[1];
};

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

struct spvDescriptorSetBuffer0
{
    constant ClearBufferConfig* config [[id(0)]];
    device Buffer* data [[id(1)]];
};

kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
    uint index = (gl_GlobalInvocationID.y * ((*spvDescriptorSet0.config).num_workgroups_x * 64u)) + gl_GlobalInvocationID.x;
    if (index < (*spvDescriptorSet0.config).buffer_bytes_div_by_four)
    {
        (*spvDescriptorSet0.data).data[index] = (*spvDescriptorSet0.config).fill_value;
    }
}


�	#

2
GLSL.std.450
main@

	
glsl\util_fill_buffer.comp��

// OpModuleProcessed entry-point main
// OpModuleProcessed client vulkan100
// OpModuleProcessed target-env vulkan1.0
// OpModuleProcessed entry-point main
#line 1
#version 450

// @[export]
// @[internal_buffer]
layout(set = 0, binding = 0) uniform ClearBufferConfig
{
    uint buffer_bytes_div_by_four;
    uint fill_value;
    uint num_workgroups_x;
} config;

layout(set = 0, binding = 1) buffer Buffer
{
    uint data[];
} data;

// Expected to invoke workgroup size where x*y >= bytes/4, and z is 1
// Keep local_size_x in sync with GROUP_SIZE_X in fill_buffer()
layout(local_size_x = 64, local_size_y = 1, local_size_z = 1) in;
void main()
{
    uint index = gl_GlobalInvocationID.y * (config.num_workgroups_x * gl_WorkGroupSize.x) + gl_GlobalInvocationID.x;
    if (index < config.buffer_bytes_div_by_four)
    {
        data.data[index] = config.fill_value;
    }
}

GL_GOOGLE_cpp_style_line_directiveGL_GOOGLE_include_directivemain	indexgl_GlobalInvocationIDClearBufferConfig
buffer_bytes_div_by_four
fill_valuenum_workgroups_xconfig)Buffer)data+dataGH#H
#H#GG"G!G(H)#G)G+"G+!
G1!  
 

;
+

 
 ; 
+ +@++!$()( *);*++-
,
1


6�;	
A
=A=��A=�>	
= 	A"!=#"�$% #�'�%&'�&
=,	A.-=/.A0+!,>0/�'�'�
8


 �
 
ClearBufferConfig


ClearBufferConfig*ClearBufferConfig.buffer_bytes_div_by_fourClearBufferConfig.fill_value"ClearBufferConfig.num_workgroups_x 

 
Buffer



data
@

main

�
 
ClearBufferConfig


ClearBufferConfig*ClearBufferConfig.buffer_bytes_div_by_fourClearBufferConfig.fill_value"ClearBufferConfig.num_workgroups_x
 

 
Buffer



data

 �
 
ClearBufferConfig


ClearBufferConfig*ClearBufferConfig.buffer_bytes_div_by_fourClearBufferConfig.fill_value"ClearBufferConfig.num_workgroups_x 

 
Buffer



data
@

main

�
 
ClearBufferConfig


ClearBufferConfig*ClearBufferConfig.buffer_bytes_div_by_fourClearBufferConfig.fill_value"ClearBufferConfig.num_workgroups_x
 

 
Buffer



data

 �
 
ClearBufferConfig


ClearBufferConfig*ClearBufferConfig.buffer_bytes_div_by_fourClearBufferConfig.fill_value"ClearBufferConfig.num_workgroups_x 

 
Buffer



data
@

main

�
 
ClearBufferConfig


ClearBufferConfig*ClearBufferConfig.buffer_bytes_div_by_fourClearBufferConfig.fill_value"ClearBufferConfig.num_workgroups_x
 

 
Buffer



data
util_fill_buffer.comp