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

Test aliased names in declared LUTs.

Browse Source
This commit is contained in:
Hans-Kristian Arntzen 2022-01-18 12:39:16 +01:00
parent 48b5a9069f
commit ac46140ba3
7 changed files with 348 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

27
reference/shaders-hlsl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp Normal file
View File

@ -0,0 +1,27 @@
static const uint3 gl_WorkGroupSize = uint3(4u, 4u, 1u);
static const int indexable[4] = { 0, 1, 2, 3 };
static const int indexable_1[4] = { 4, 5, 6, 7 };
RWByteAddressBuffer _6 : register(u0);
static uint3 gl_LocalInvocationID;
static uint3 gl_GlobalInvocationID;
struct SPIRV_Cross_Input
{
uint3 gl_LocalInvocationID : SV_GroupThreadID;
uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
};
void comp_main()
{
_6.Store(gl_GlobalInvocationID.x * 4 + 0, uint(indexable[gl_LocalInvocationID.x] + indexable_1[gl_LocalInvocationID.y]));
}
[numthreads(4, 4, 1)]
void main(SPIRV_Cross_Input stage_input)
{
gl_LocalInvocationID = stage_input.gl_LocalInvocationID;
gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID;
comp_main();
}

61
reference/shaders-msl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp Normal file
View File

@ -0,0 +1,61 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"
#pragma clang diagnostic ignored "-Wmissing-braces"
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
template<typename T, size_t Num>
struct spvUnsafeArray
{
T elements[Num ? Num : 1];
thread T& operator [] (size_t pos) thread
{
return elements[pos];
}
constexpr const thread T& operator [] (size_t pos) const thread
{
return elements[pos];
}
device T& operator [] (size_t pos) device
{
return elements[pos];
}
constexpr const device T& operator [] (size_t pos) const device
{
return elements[pos];
}
constexpr const constant T& operator [] (size_t pos) const constant
{
return elements[pos];
}
threadgroup T& operator [] (size_t pos) threadgroup
{
return elements[pos];
}
constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
{
return elements[pos];
}
};
struct SSBO
{
int values[1];
};
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 4u, 1u);
constant spvUnsafeArray<int, 4> indexable = spvUnsafeArray<int, 4>({ 0, 1, 2, 3 });
constant spvUnsafeArray<int, 4> indexable_1 = spvUnsafeArray<int, 4>({ 4, 5, 6, 7 });
kernel void main0(device SSBO& _6 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
{
_6.values[gl_GlobalInvocationID.x] = indexable[gl_LocalInvocationID.x] + indexable_1[gl_LocalInvocationID.y];
}

16
reference/shaders-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp Normal file
View File

@ -0,0 +1,16 @@
#version 450
layout(local_size_x = 4, local_size_y = 4, local_size_z = 1) in;
const int indexable[4] = int[](0, 1, 2, 3);
const int indexable_1[4] = int[](4, 5, 6, 7);
layout(binding = 0, std430) buffer SSBO
{
int values[];
} _6;
void main()
{
_6.values[gl_GlobalInvocationID.x] = indexable[gl_LocalInvocationID.x] + indexable_1[gl_LocalInvocationID.y];
}

81
shaders-hlsl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp Normal file
View File

@ -0,0 +1,81 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 10
; Bound: 49
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID %gl_LocalInvocationID
OpExecutionMode %main LocalSize 4 4 1
OpSource GLSL 450
OpName %main "main"
OpName %SSBO "SSBO"
OpMemberName %SSBO 0 "values"
OpName %_ ""
OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
OpName %gl_LocalInvocationID "gl_LocalInvocationID"
OpName %indexable "indexable"
OpName %indexable_0 "indexable"
OpName %25 "indexable"
OpName %38 "indexable"
OpDecorate %_runtimearr_int ArrayStride 4
OpMemberDecorate %SSBO 0 Offset 0
OpDecorate %SSBO BufferBlock
OpDecorate %_ DescriptorSet 0
OpDecorate %_ Binding 0
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
%void = OpTypeVoid
%3 = OpTypeFunction %void
%int = OpTypeInt 32 1
%_runtimearr_int = OpTypeRuntimeArray %int
%SSBO = OpTypeStruct %_runtimearr_int
%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
%_ = OpVariable %_ptr_Uniform_SSBO Uniform
%int_0 = OpConstant %int 0
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
%uint_0 = OpConstant %uint 0
%_ptr_Input_uint = OpTypePointer Input %uint
%uint_4 = OpConstant %uint 4
%_arr_int_uint_4 = OpTypeArray %int %uint_4
%int_1 = OpConstant %int 1
%int_2 = OpConstant %int 2
%int_3 = OpConstant %int 3
%25 = OpConstantComposite %_arr_int_uint_4 %int_0 %int_1 %int_2 %int_3
%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input
%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4
%_ptr_Function_int = OpTypePointer Function %int
%int_4 = OpConstant %int 4
%int_5 = OpConstant %int 5
%int_6 = OpConstant %int 6
%int_7 = OpConstant %int 7
%38 = OpConstantComposite %_arr_int_uint_4 %int_4 %int_5 %int_6 %int_7
%uint_1 = OpConstant %uint 1
%_ptr_Uniform_int = OpTypePointer Uniform %int
%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_4 %uint_4 %uint_1
%main = OpFunction %void None %3
%5 = OpLabel
%indexable = OpVariable %_ptr_Function__arr_int_uint_4 Function
%indexable_0 = OpVariable %_ptr_Function__arr_int_uint_4 Function
%18 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
%19 = OpLoad %uint %18
%27 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0
%28 = OpLoad %uint %27
OpStore %indexable %25
%32 = OpAccessChain %_ptr_Function_int %indexable %28
%33 = OpLoad %int %32
%40 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_1
%41 = OpLoad %uint %40
OpStore %indexable_0 %38
%43 = OpAccessChain %_ptr_Function_int %indexable_0 %41
%44 = OpLoad %int %43
%45 = OpIAdd %int %33 %44
%47 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %19
OpStore %47 %45
OpReturn
OpFunctionEnd

81
shaders-msl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp Normal file
View File

@ -0,0 +1,81 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 10
; Bound: 49
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID %gl_LocalInvocationID
OpExecutionMode %main LocalSize 4 4 1
OpSource GLSL 450
OpName %main "main"
OpName %SSBO "SSBO"
OpMemberName %SSBO 0 "values"
OpName %_ ""
OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
OpName %gl_LocalInvocationID "gl_LocalInvocationID"
OpName %indexable "indexable"
OpName %indexable_0 "indexable"
OpName %25 "indexable"
OpName %38 "indexable"
OpDecorate %_runtimearr_int ArrayStride 4
OpMemberDecorate %SSBO 0 Offset 0
OpDecorate %SSBO BufferBlock
OpDecorate %_ DescriptorSet 0
OpDecorate %_ Binding 0
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
%void = OpTypeVoid
%3 = OpTypeFunction %void
%int = OpTypeInt 32 1
%_runtimearr_int = OpTypeRuntimeArray %int
%SSBO = OpTypeStruct %_runtimearr_int
%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
%_ = OpVariable %_ptr_Uniform_SSBO Uniform
%int_0 = OpConstant %int 0
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
%uint_0 = OpConstant %uint 0
%_ptr_Input_uint = OpTypePointer Input %uint
%uint_4 = OpConstant %uint 4
%_arr_int_uint_4 = OpTypeArray %int %uint_4
%int_1 = OpConstant %int 1
%int_2 = OpConstant %int 2
%int_3 = OpConstant %int 3
%25 = OpConstantComposite %_arr_int_uint_4 %int_0 %int_1 %int_2 %int_3
%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input
%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4
%_ptr_Function_int = OpTypePointer Function %int
%int_4 = OpConstant %int 4
%int_5 = OpConstant %int 5
%int_6 = OpConstant %int 6
%int_7 = OpConstant %int 7
%38 = OpConstantComposite %_arr_int_uint_4 %int_4 %int_5 %int_6 %int_7
%uint_1 = OpConstant %uint 1
%_ptr_Uniform_int = OpTypePointer Uniform %int
%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_4 %uint_4 %uint_1
%main = OpFunction %void None %3
%5 = OpLabel
%indexable = OpVariable %_ptr_Function__arr_int_uint_4 Function
%indexable_0 = OpVariable %_ptr_Function__arr_int_uint_4 Function
%18 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
%19 = OpLoad %uint %18
%27 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0
%28 = OpLoad %uint %27
OpStore %indexable %25
%32 = OpAccessChain %_ptr_Function_int %indexable %28
%33 = OpLoad %int %32
%40 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_1
%41 = OpLoad %uint %40
OpStore %indexable_0 %38
%43 = OpAccessChain %_ptr_Function_int %indexable_0 %41
%44 = OpLoad %int %43
%45 = OpIAdd %int %33 %44
%47 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %19
OpStore %47 %45
OpReturn
OpFunctionEnd

81
shaders-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp Normal file
View File

@ -0,0 +1,81 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 10
; Bound: 49
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID %gl_LocalInvocationID
OpExecutionMode %main LocalSize 4 4 1
OpSource GLSL 450
OpName %main "main"
OpName %SSBO "SSBO"
OpMemberName %SSBO 0 "values"
OpName %_ ""
OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
OpName %gl_LocalInvocationID "gl_LocalInvocationID"
OpName %indexable "indexable"
OpName %indexable_0 "indexable"
OpName %25 "indexable"
OpName %38 "indexable"
OpDecorate %_runtimearr_int ArrayStride 4
OpMemberDecorate %SSBO 0 Offset 0
OpDecorate %SSBO BufferBlock
OpDecorate %_ DescriptorSet 0
OpDecorate %_ Binding 0
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
%void = OpTypeVoid
%3 = OpTypeFunction %void
%int = OpTypeInt 32 1
%_runtimearr_int = OpTypeRuntimeArray %int
%SSBO = OpTypeStruct %_runtimearr_int
%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
%_ = OpVariable %_ptr_Uniform_SSBO Uniform
%int_0 = OpConstant %int 0
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
%uint_0 = OpConstant %uint 0
%_ptr_Input_uint = OpTypePointer Input %uint
%uint_4 = OpConstant %uint 4
%_arr_int_uint_4 = OpTypeArray %int %uint_4
%int_1 = OpConstant %int 1
%int_2 = OpConstant %int 2
%int_3 = OpConstant %int 3
%25 = OpConstantComposite %_arr_int_uint_4 %int_0 %int_1 %int_2 %int_3
%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input
%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4
%_ptr_Function_int = OpTypePointer Function %int
%int_4 = OpConstant %int 4
%int_5 = OpConstant %int 5
%int_6 = OpConstant %int 6
%int_7 = OpConstant %int 7
%38 = OpConstantComposite %_arr_int_uint_4 %int_4 %int_5 %int_6 %int_7
%uint_1 = OpConstant %uint 1
%_ptr_Uniform_int = OpTypePointer Uniform %int
%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_4 %uint_4 %uint_1
%main = OpFunction %void None %3
%5 = OpLabel
%indexable = OpVariable %_ptr_Function__arr_int_uint_4 Function
%indexable_0 = OpVariable %_ptr_Function__arr_int_uint_4 Function
%18 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
%19 = OpLoad %uint %18
%27 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0
%28 = OpLoad %uint %27
OpStore %indexable %25
%32 = OpAccessChain %_ptr_Function_int %indexable %28
%33 = OpLoad %int %32
%40 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_1
%41 = OpLoad %uint %40
OpStore %indexable_0 %38
%43 = OpAccessChain %_ptr_Function_int %indexable_0 %41
%44 = OpLoad %int %43
%45 = OpIAdd %int %33 %44
%47 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %19
OpStore %47 %45
OpReturn
OpFunctionEnd

1
spirv_hlsl.cpp
View File

@ -1167,6 +1167,7 @@ void CompilerHLSL::emit_composite_constants()
if (type.basetype == SPIRType::Struct || !type.array.empty())
{
add_resource_name(c.self);
auto name = to_name(c.self);
statement("static const ", variable_decl(type, name), " = ", constant_expression(c), ";");
emitted = true;