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

Test loading from and storing to packed vectors.

Browse Source
This commit is contained in:
Chip Davis 2018-11-14 10:47:20 -06:00
parent bed4918cb5
commit a5882da091
3 changed files with 48 additions and 8 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
reference/opt/shaders-msl/asm/comp/relaxed-block-layout.asm.comp
View File

@ -12,9 +12,11 @@ struct foo
packed_half2 wibble;
};
kernel void main0(device foo& _6 [[buffer(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
kernel void main0(device foo& _8 [[buffer(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_NumWorkGroups [[threadgroups_per_grid]])
{
_6.bar = gl_LocalInvocationID.x;
_6.baz = float3(gl_GlobalInvocationID);
_8.bar = gl_LocalInvocationID.x;
_8.baz = float3(gl_GlobalInvocationID);
_8.blah = uchar4(uint4(uint4(uchar4(_8.blah)).xyz + gl_WorkGroupID, 0u));
_8.wibble = half2(float2(half2(_8.wibble)) * float2(gl_NumWorkGroups.xy));
}

8
reference/shaders-msl/asm/comp/relaxed-block-layout.asm.comp
View File

@ -12,9 +12,11 @@ struct foo
packed_half2 wibble;
};
kernel void main0(device foo& _6 [[buffer(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
kernel void main0(device foo& _8 [[buffer(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_NumWorkGroups [[threadgroups_per_grid]])
{
_6.bar = gl_LocalInvocationID.x;
_6.baz = float3(gl_GlobalInvocationID);
_8.bar = gl_LocalInvocationID.x;
_8.baz = float3(gl_GlobalInvocationID);
_8.blah = uchar4(uint4(uint4(uchar4(_8.blah)).xyz + gl_WorkGroupID, 0u));
_8.wibble = half2(float2(half2(_8.wibble)) * float2(gl_NumWorkGroups.xy));
}

40
shaders-msl/asm/comp/relaxed-block-layout.asm.comp
View File

@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Khronos Glslang Reference Front End; 7
; Bound: 33
; Bound: 63
; Schema: 0
OpCapability Shader
OpCapability StorageBuffer16BitAccess
@ -10,7 +10,7 @@
OpExtension "SPV_KHR_8bit_storage"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %gl_LocalInvocationID %gl_GlobalInvocationID
OpEntryPoint GLCompute %main "main" %gl_LocalInvocationID %gl_GlobalInvocationID %gl_WorkGroupID %gl_NumWorkGroups
OpExecutionMode %main LocalSize 1 1 1
OpSource GLSL 450
OpSourceExtension "GL_EXT_shader_16bit_storage"
@ -25,6 +25,8 @@
OpName %_ ""
OpName %gl_LocalInvocationID "gl_LocalInvocationID"
OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
OpName %gl_WorkGroupID "gl_WorkGroupID"
OpName %gl_NumWorkGroups "gl_NumWorkGroups"
OpMemberDecorate %foo 0 Offset 0
OpMemberDecorate %foo 1 Offset 4
OpMemberDecorate %foo 2 Offset 16
@ -35,6 +37,8 @@
OpDecorate %_ Binding 0
OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
OpDecorate %gl_WorkGroupID BuiltIn WorkgroupId
OpDecorate %gl_NumWorkGroups BuiltIn NumWorkgroups
%void = OpTypeVoid
%3 = OpTypeFunction %void
%uint = OpTypeInt 32 0
@ -58,6 +62,15 @@
%int_1 = OpConstant %int 1
%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
%int_3 = OpConstant %int 3
%_ptr_Uniform_v4uchar = OpTypePointer Uniform %v4uchar
%v4uint = OpTypeVector %uint 4
%gl_WorkGroupID = OpVariable %_ptr_Input_v3uint Input
%int_4 = OpConstant %int 4
%_ptr_Uniform_v2half = OpTypePointer Uniform %v2half
%v2float = OpTypeVector %float 2
%gl_NumWorkGroups = OpVariable %_ptr_Input_v3uint Input
%v2uint = OpTypeVector %uint 2
%main = OpFunction %void None %3
%5 = OpLabel
%23 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0
@ -68,5 +81,28 @@
%30 = OpConvertUToF %v3float %29
%32 = OpAccessChain %_ptr_Uniform_v3float %_ %int_1
OpStore %32 %30
%35 = OpAccessChain %_ptr_Uniform_v4uchar %_ %int_3
%36 = OpLoad %v4uchar %35
%38 = OpUConvert %v4uint %36
%39 = OpVectorShuffle %v3uint %38 %38 0 1 2
%41 = OpLoad %v3uint %gl_WorkGroupID
%42 = OpIAdd %v3uint %39 %41
%43 = OpCompositeExtract %uint %42 0
%44 = OpCompositeExtract %uint %42 1
%45 = OpCompositeExtract %uint %42 2
%46 = OpCompositeConstruct %v4uint %43 %44 %45 %uint_0
%47 = OpUConvert %v4uchar %46
%48 = OpAccessChain %_ptr_Uniform_v4uchar %_ %int_3
OpStore %48 %47
%51 = OpAccessChain %_ptr_Uniform_v2half %_ %int_4
%52 = OpLoad %v2half %51
%54 = OpFConvert %v2float %52
%57 = OpLoad %v3uint %gl_NumWorkGroups
%58 = OpVectorShuffle %v2uint %57 %57 0 1
%59 = OpConvertUToF %v2float %58
%60 = OpFMul %v2float %54 %59
%61 = OpFConvert %v2half %60
%62 = OpAccessChain %_ptr_Uniform_v2half %_ %int_4
OpStore %62 %61
OpReturn
OpFunctionEnd