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

Merge pull request #912 from KhronosGroup/fix-910

Browse Source 
Parser: Fix OpCompositeConstruct with OpUndef.
This commit is contained in:
Hans-Kristian Arntzen 2019-03-27 12:20:41 +01:00 committed by GitHub
commit b64e4564ba
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
7 changed files with 174 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

14
reference/shaders-hlsl-no-opt/asm/comp/constant-composite-undef.asm.comp Normal file
View File

@ -0,0 +1,14 @@
RWByteAddressBuffer block : register(u0);
float _15;
void comp_main()
{
block.Store4(0, asuint(float4(0.100000001490116119384765625f, 0.20000000298023223876953125f, 0.300000011920928955078125f, 0.0f)));
}
[numthreads(1, 1, 1)]
void main()
{
comp_main();
}

17
reference/shaders-msl-no-opt/asm/comp/constant-composite-undef.asm.comp Normal file
View File

@ -0,0 +1,17 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct Block
{
float4 f;
};
constant float _15 = {};
kernel void main0(device Block& block [[buffer(0)]])
{
block.f = float4(0.100000001490116119384765625, 0.20000000298023223876953125, 0.300000011920928955078125, 0.0);
}

15
reference/shaders-no-opt/asm/comp/constant-composite-undef.asm.comp Normal file
View File

@ -0,0 +1,15 @@
#version 450
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(binding = 0, std430) buffer Block
{
vec4 f;
} block;
float _15;
void main()
{
block.f = vec4(0.100000001490116119384765625, 0.20000000298023223876953125, 0.300000011920928955078125, 0.0);
}

40
shaders-hlsl-no-opt/asm/comp/constant-composite-undef.asm.comp Normal file
View File

@ -0,0 +1,40 @@
; SPIR-V
; Version: 1.3
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 20
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpSource GLSL 450
OpName %main "main"
OpName %Block "Block"
OpMemberName %Block 0 "f"
OpName %block "block"
OpMemberDecorate %Block 0 Offset 0
OpDecorate %Block BufferBlock
OpDecorate %block DescriptorSet 0
OpDecorate %block Binding 0
%void = OpTypeVoid
%6 = OpTypeFunction %void
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%Block = OpTypeStruct %v4float
%_ptr_Uniform_Block = OpTypePointer Uniform %Block
%block = OpVariable %_ptr_Uniform_Block Uniform
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%float_0_100000001 = OpConstant %float 0.100000001
%float_0_200000003 = OpConstant %float 0.200000003
%float_0_300000012 = OpConstant %float 0.300000012
%15 = OpUndef %float
%16 = OpConstantComposite %v4float %float_0_100000001 %float_0_200000003 %float_0_300000012 %15
%_ptr_Uniform_v4float = OpTypePointer Uniform %v4float
%main = OpFunction %void None %6
%18 = OpLabel
%19 = OpAccessChain %_ptr_Uniform_v4float %block %int_0
OpStore %19 %16
OpReturn
OpFunctionEnd

40
shaders-msl-no-opt/asm/comp/constant-composite-undef.asm.comp Normal file
View File

@ -0,0 +1,40 @@
; SPIR-V
; Version: 1.3
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 20
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpSource GLSL 450
OpName %main "main"
OpName %Block "Block"
OpMemberName %Block 0 "f"
OpName %block "block"
OpMemberDecorate %Block 0 Offset 0
OpDecorate %Block BufferBlock
OpDecorate %block DescriptorSet 0
OpDecorate %block Binding 0
%void = OpTypeVoid
%6 = OpTypeFunction %void
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%Block = OpTypeStruct %v4float
%_ptr_Uniform_Block = OpTypePointer Uniform %Block
%block = OpVariable %_ptr_Uniform_Block Uniform
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%float_0_100000001 = OpConstant %float 0.100000001
%float_0_200000003 = OpConstant %float 0.200000003
%float_0_300000012 = OpConstant %float 0.300000012
%15 = OpUndef %float
%16 = OpConstantComposite %v4float %float_0_100000001 %float_0_200000003 %float_0_300000012 %15
%_ptr_Uniform_v4float = OpTypePointer Uniform %v4float
%main = OpFunction %void None %6
%18 = OpLabel
%19 = OpAccessChain %_ptr_Uniform_v4float %block %int_0
OpStore %19 %16
OpReturn
OpFunctionEnd

40
shaders-no-opt/asm/comp/constant-composite-undef.asm.comp Normal file
View File

@ -0,0 +1,40 @@
; SPIR-V
; Version: 1.3
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 20
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpSource GLSL 450
OpName %main "main"
OpName %Block "Block"
OpMemberName %Block 0 "f"
OpName %block "block"
OpMemberDecorate %Block 0 Offset 0
OpDecorate %Block BufferBlock
OpDecorate %block DescriptorSet 0
OpDecorate %block Binding 0
%void = OpTypeVoid
%6 = OpTypeFunction %void
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%Block = OpTypeStruct %v4float
%_ptr_Uniform_Block = OpTypePointer Uniform %Block
%block = OpVariable %_ptr_Uniform_Block Uniform
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%float_0_100000001 = OpConstant %float 0.100000001
%float_0_200000003 = OpConstant %float 0.200000003
%float_0_300000012 = OpConstant %float 0.300000012
%15 = OpUndef %float
%16 = OpConstantComposite %v4float %float_0_100000001 %float_0_200000003 %float_0_300000012 %15
%_ptr_Uniform_v4float = OpTypePointer Uniform %v4float
%main = OpFunction %void None %6
%18 = OpLabel
%19 = OpAccessChain %_ptr_Uniform_v4float %block %int_0
OpStore %19 %16
OpReturn
OpFunctionEnd

8
spirv_parser.cpp
View File

@ -775,6 +775,7 @@ void Parser::parse(const Instruction &instruction)
// We do not know their value, so any attempt to query SPIRConstant later
// will fail. We can only propagate the ID of the expression and use to_expression on it.
auto *constant_op = maybe_get<SPIRConstantOp>(ops[2 + i]);
auto *undef_op = maybe_get<SPIRUndef>(ops[2 + i]);
if (constant_op)
{
if (op == OpConstantComposite)
@ -786,6 +787,13 @@ void Parser::parse(const Instruction &instruction)
remapped_constant_ops[i].specialization = true;
c[i] = &remapped_constant_ops[i];
}
else if (undef_op)
{
// Undefined, just pick 0.
remapped_constant_ops[i].make_null(get<SPIRType>(undef_op->basetype));
remapped_constant_ops[i].constant_type = undef_op->basetype;
c[i] = &remapped_constant_ops[i];
}
else
c[i] = &get<SPIRConstant>(ops[2 + i]);
}