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

Add support for LocalSizeId.

Browse Source 
WorkgroupSize builtin is deprecated in 1.6 and LocalSizeId is supported
in Vulkan starting with maintenance4.
This commit is contained in:
Hans-Kristian Arntzen 2022-01-05 15:53:51 +01:00
parent 35bb328443
commit 7c83fc22fa
23 changed files with 800 additions and 37 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
reference/opt/shaders-hlsl/asm/vert/spec-constant-op-composite.asm.vert
View File

@ -8,10 +8,7 @@ static const int _20 = (_7 + 2);
#endif
static const uint _8 = SPIRV_CROSS_CONSTANT_ID_202;
static const uint _25 = (_8 % 5u);
#ifndef SPIRV_CROSS_CONSTANT_ID_0
#define SPIRV_CROSS_CONSTANT_ID_0 int4(20, 30, _20, _20)
#endif
static const int4 _30 = SPIRV_CROSS_CONSTANT_ID_0;
static const int4 _30 = int4(20, 30, _20, _20);
static const int2 _32 = int2(_30.y, _30.x);
static const int _33 = _30.y;

37
reference/shaders-hlsl-no-opt/asm/comp/local-size-id-override.asm.comp Normal file
View File

@ -0,0 +1,37 @@
#ifndef SPIRV_CROSS_CONSTANT_ID_1
#define SPIRV_CROSS_CONSTANT_ID_1 11u
#endif
static const uint _10 = SPIRV_CROSS_CONSTANT_ID_1;
#ifndef SPIRV_CROSS_CONSTANT_ID_2
#define SPIRV_CROSS_CONSTANT_ID_2 12u
#endif
static const uint _11 = SPIRV_CROSS_CONSTANT_ID_2;
#ifndef SPIRV_CROSS_CONSTANT_ID_3
#define SPIRV_CROSS_CONSTANT_ID_3 13u
#endif
static const uint _4 = SPIRV_CROSS_CONSTANT_ID_3;
#ifndef SPIRV_CROSS_CONSTANT_ID_4
#define SPIRV_CROSS_CONSTANT_ID_4 14u
#endif
static const uint _5 = SPIRV_CROSS_CONSTANT_ID_4;
static const uint3 gl_WorkGroupSize = uint3(3u, _10, _11);
RWByteAddressBuffer _8 : register(u0);
static uint3 gl_GlobalInvocationID;
struct SPIRV_Cross_Input
{
uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
};
void comp_main()
{
_8.Store4(gl_GlobalInvocationID.x * 16 + 0, asuint(asfloat(_8.Load4(gl_GlobalInvocationID.x * 16 + 0)) + 2.0f.xxxx));
}
[numthreads(3, SPIRV_CROSS_CONSTANT_ID_1, SPIRV_CROSS_CONSTANT_ID_2)]
void main(SPIRV_Cross_Input stage_input)
{
gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID;
comp_main();
}

38
reference/shaders-hlsl-no-opt/asm/comp/local-size-id.asm.comp Normal file
View File

@ -0,0 +1,38 @@
#ifndef SPIRV_CROSS_CONSTANT_ID_1
#define SPIRV_CROSS_CONSTANT_ID_1 11
#endif
static const int _10 = SPIRV_CROSS_CONSTANT_ID_1;
#ifndef SPIRV_CROSS_CONSTANT_ID_2
#define SPIRV_CROSS_CONSTANT_ID_2 12
#endif
static const int _11 = SPIRV_CROSS_CONSTANT_ID_2;
#ifndef SPIRV_CROSS_CONSTANT_ID_3
#define SPIRV_CROSS_CONSTANT_ID_3 13
#endif
static const int _4 = SPIRV_CROSS_CONSTANT_ID_3;
#ifndef SPIRV_CROSS_CONSTANT_ID_4
#define SPIRV_CROSS_CONSTANT_ID_4 14
#endif
static const int _5 = SPIRV_CROSS_CONSTANT_ID_4;
static const uint _29 = (uint(_4) + 3u);
static const uint3 _30 = uint3(_29, _5, 2u);
RWByteAddressBuffer _8 : register(u0);
static uint3 gl_GlobalInvocationID;
struct SPIRV_Cross_Input
{
uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
};
void comp_main()
{
_8.Store4(gl_GlobalInvocationID.x * 16 + 0, asuint(((((asfloat(_8.Load4(gl_GlobalInvocationID.x * 16 + 0)) + 2.0f.xxxx) + float3(_30).xyzz) * float(_4)) * float(_5)) * float(int(2u))));
}
[numthreads(SPIRV_CROSS_CONSTANT_ID_3, SPIRV_CROSS_CONSTANT_ID_4, 2)]
void main(SPIRV_Cross_Input stage_input)
{
gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID;
comp_main();
}

5
reference/shaders-hlsl/asm/vert/spec-constant-op-composite.asm.vert
View File

@ -8,10 +8,7 @@ static const int _20 = (_7 + 2);
#endif
static const uint _8 = SPIRV_CROSS_CONSTANT_ID_202;
static const uint _25 = (_8 % 5u);
#ifndef SPIRV_CROSS_CONSTANT_ID_0
#define SPIRV_CROSS_CONSTANT_ID_0 int4(20, 30, _20, _20)
#endif
static const int4 _30 = SPIRV_CROSS_CONSTANT_ID_0;
static const int4 _30 = int4(20, 30, _20, _20);
static const int2 _32 = int2(_30.y, _30.x);
static const int _33 = _30.y;
#ifndef SPIRV_CROSS_CONSTANT_ID_200

25
reference/shaders-msl-no-opt/asm/comp/local-size-id-override.asm.comp Normal file
View File

@ -0,0 +1,25 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct SSBO
{
float4 values[1];
};
constant uint _10_tmp [[function_constant(1)]];
constant uint _10 = is_function_constant_defined(_10_tmp) ? _10_tmp : 11u;
constant uint _11_tmp [[function_constant(2)]];
constant uint _11 = is_function_constant_defined(_11_tmp) ? _11_tmp : 12u;
constant uint _4_tmp [[function_constant(3)]];
constant uint _4 = is_function_constant_defined(_4_tmp) ? _4_tmp : 13u;
constant uint _5_tmp [[function_constant(4)]];
constant uint _5 = is_function_constant_defined(_5_tmp) ? _5_tmp : 14u;
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(3u, _10, _11);
kernel void main0(device SSBO& _8 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
_8.values[gl_GlobalInvocationID.x] += float4(2.0);
}

26
reference/shaders-msl-no-opt/asm/comp/local-size-id.asm.comp Normal file
View File

@ -0,0 +1,26 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct SSBO
{
float4 values[1];
};
constant int _10_tmp [[function_constant(1)]];
constant int _10 = is_function_constant_defined(_10_tmp) ? _10_tmp : 11;
constant int _11_tmp [[function_constant(2)]];
constant int _11 = is_function_constant_defined(_11_tmp) ? _11_tmp : 12;
constant int _4_tmp [[function_constant(3)]];
constant int _4 = is_function_constant_defined(_4_tmp) ? _4_tmp : 13;
constant int _5_tmp [[function_constant(4)]];
constant int _5 = is_function_constant_defined(_5_tmp) ? _5_tmp : 14;
constant uint _29 = (uint(_4) + 3u);
constant uint3 _30 = uint3(_29, _5, 2u);
kernel void main0(device SSBO& _8 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
_8.values[gl_GlobalInvocationID.x] = ((((_8.values[gl_GlobalInvocationID.x] + float4(2.0)) + float3(_30).xyzz) * float(_4)) * float(_5)) * float(int(2u));
}

29
reference/shaders-no-opt/asm/comp/local-size-id-override.vk.asm.comp Normal file
View File

@ -0,0 +1,29 @@
#version 450
#ifndef SPIRV_CROSS_CONSTANT_ID_1
#define SPIRV_CROSS_CONSTANT_ID_1 11u
#endif
#ifndef SPIRV_CROSS_CONSTANT_ID_2
#define SPIRV_CROSS_CONSTANT_ID_2 12u
#endif
#ifndef SPIRV_CROSS_CONSTANT_ID_3
#define SPIRV_CROSS_CONSTANT_ID_3 13u
#endif
const uint _4 = SPIRV_CROSS_CONSTANT_ID_3;
#ifndef SPIRV_CROSS_CONSTANT_ID_4
#define SPIRV_CROSS_CONSTANT_ID_4 14u
#endif
const uint _5 = SPIRV_CROSS_CONSTANT_ID_4;
layout(local_size_x = 3, local_size_y = SPIRV_CROSS_CONSTANT_ID_1, local_size_z = SPIRV_CROSS_CONSTANT_ID_2) in;
layout(binding = 0, std430) buffer SSBO
{
vec4 values[];
} _8;
void main()
{
_8.values[gl_GlobalInvocationID.x] += vec4(2.0);
}

16
reference/shaders-no-opt/asm/comp/local-size-id-override.vk.asm.comp.vk Normal file
View File

@ -0,0 +1,16 @@
#version 450
layout(local_size_x = 3, local_size_y_id = 1, local_size_z_id = 2) in;
layout(constant_id = 3) const uint _4 = 13u;
layout(constant_id = 4) const uint _5 = 14u;
layout(set = 0, binding = 0, std430) buffer SSBO
{
vec4 values[];
} _8;
void main()
{
_8.values[gl_GlobalInvocationID.x] += vec4(2.0);
}

31
reference/shaders-no-opt/asm/comp/local-size-id.vk.asm.comp Normal file
View File

@ -0,0 +1,31 @@
#version 450
#ifndef SPIRV_CROSS_CONSTANT_ID_1
#define SPIRV_CROSS_CONSTANT_ID_1 11
#endif
const int _10 = SPIRV_CROSS_CONSTANT_ID_1;
#ifndef SPIRV_CROSS_CONSTANT_ID_2
#define SPIRV_CROSS_CONSTANT_ID_2 12
#endif
const int _11 = SPIRV_CROSS_CONSTANT_ID_2;
#ifndef SPIRV_CROSS_CONSTANT_ID_3
#define SPIRV_CROSS_CONSTANT_ID_3 13
#endif
#ifndef SPIRV_CROSS_CONSTANT_ID_4
#define SPIRV_CROSS_CONSTANT_ID_4 14
#endif
const uint _29 = (uint(int(gl_WorkGroupSize.x)) + 3u);
const uvec3 _30 = uvec3(_29, int(gl_WorkGroupSize.y), 2u);
layout(local_size_x = SPIRV_CROSS_CONSTANT_ID_3, local_size_y = SPIRV_CROSS_CONSTANT_ID_4, local_size_z = 2) in;
layout(binding = 0, std430) buffer SSBO
{
vec4 values[];
} _8;
void main()
{
_8.values[gl_GlobalInvocationID.x] = ((((_8.values[gl_GlobalInvocationID.x] + vec4(2.0)) + vec3(_30).xyzz) * float(int(gl_WorkGroupSize.x))) * float(int(gl_WorkGroupSize.y))) * float(int(2u));
}

18
reference/shaders-no-opt/asm/comp/local-size-id.vk.asm.comp.vk Normal file
View File

@ -0,0 +1,18 @@
#version 450
layout(local_size_x_id = 3, local_size_y_id = 4, local_size_z = 2) in;
layout(constant_id = 1) const int _10 = 11;
layout(constant_id = 2) const int _11 = 12;
const uint _29 = (uint(int(gl_WorkGroupSize.x)) + 3u);
const uvec3 _30 = uvec3(_29, int(gl_WorkGroupSize.y), 2u);
layout(set = 0, binding = 0, std430) buffer SSBO
{
vec4 values[];
} _8;
void main()
{
_8.values[gl_GlobalInvocationID.x] = ((((_8.values[gl_GlobalInvocationID.x] + vec4(2.0)) + vec3(_30).xyzz) * float(int(gl_WorkGroupSize.x))) * float(int(gl_WorkGroupSize.y))) * float(int(2u));
}

60
shaders-hlsl-no-opt/asm/comp/local-size-id-override.asm.comp Normal file
View File

@ -0,0 +1,60 @@
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID
OpExecutionModeId %main LocalSizeId %spec_3 %spec_4 %uint_2
OpSource GLSL 450
OpName %main "main"
OpName %SSBO "SSBO"
OpMemberName %SSBO 0 "values"
OpName %_ ""
OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
OpDecorate %_runtimearr_v4float ArrayStride 16
OpMemberDecorate %SSBO 0 Offset 0
OpDecorate %SSBO Block
OpDecorate %_ DescriptorSet 0
OpDecorate %_ Binding 0
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
OpDecorate %spec_1 SpecId 1
OpDecorate %spec_2 SpecId 2
OpDecorate %spec_3 SpecId 3
OpDecorate %spec_4 SpecId 4
OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
%void = OpTypeVoid
%3 = OpTypeFunction %void
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_runtimearr_v4float = OpTypeRuntimeArray %v4float
%SSBO = OpTypeStruct %_runtimearr_v4float
%_ptr_Uniform_SSBO = OpTypePointer StorageBuffer %SSBO
%_ = OpVariable %_ptr_Uniform_SSBO StorageBuffer
%int = OpTypeInt 32 1
%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
%float_2 = OpConstant %float 2
%_ptr_Uniform_v4float = OpTypePointer StorageBuffer %v4float
%spec_1 = OpSpecConstant %uint 11
%spec_2 = OpSpecConstant %uint 12
%spec_3 = OpSpecConstant %uint 13
%spec_4 = OpSpecConstant %uint 14
%uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%uint_3 = OpConstant %uint 3
%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %uint_3 %spec_1 %spec_2
%main = OpFunction %void None %3
%5 = OpLabel
%20 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
%21 = OpLoad %uint %20
%24 = OpAccessChain %_ptr_Uniform_v4float %_ %int_0 %21
%25 = OpLoad %v4float %24
%26 = OpCompositeConstruct %v4float %float_2 %float_2 %float_2 %float_2
%27 = OpFAdd %v4float %25 %26
%28 = OpAccessChain %_ptr_Uniform_v4float %_ %int_0 %21
OpStore %28 %27
OpReturn
OpFunctionEnd

76
shaders-hlsl-no-opt/asm/comp/local-size-id.asm.comp Normal file
View File

@ -0,0 +1,76 @@
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID
OpExecutionModeId %main LocalSizeId %spec_3 %spec_4 %uint_2
OpSource GLSL 450
OpName %main "main"
OpName %SSBO "SSBO"
OpMemberName %SSBO 0 "values"
OpName %_ ""
OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
OpDecorate %_runtimearr_v4float ArrayStride 16
OpMemberDecorate %SSBO 0 Offset 0
OpDecorate %SSBO Block
OpDecorate %_ DescriptorSet 0
OpDecorate %_ Binding 0
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
OpDecorate %spec_1 SpecId 1
OpDecorate %spec_2 SpecId 2
OpDecorate %spec_3 SpecId 3
OpDecorate %spec_4 SpecId 4
%void = OpTypeVoid
%3 = OpTypeFunction %void
%float = OpTypeFloat 32
%v3float = OpTypeVector %float 3
%v4float = OpTypeVector %float 4
%_runtimearr_v4float = OpTypeRuntimeArray %v4float
%SSBO = OpTypeStruct %_runtimearr_v4float
%_ptr_Uniform_SSBO = OpTypePointer StorageBuffer %SSBO
%_ = OpVariable %_ptr_Uniform_SSBO StorageBuffer
%int = OpTypeInt 32 1
%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
%float_2 = OpConstant %float 2
%_ptr_Uniform_v4float = OpTypePointer StorageBuffer %v4float
; Test that we can declare the spec constant as signed.
; Needs implicit bitcast since WorkGroupSize is uint.
%spec_1 = OpSpecConstant %int 11
%spec_2 = OpSpecConstant %int 12
%spec_3 = OpSpecConstant %int 13
%spec_4 = OpSpecConstant %int 14
%uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%uint_3 = OpConstant %uint 3
; Test that we can build spec constant composites out of local size id values.
; Needs special case handling.
%spec_3_op = OpSpecConstantOp %uint IAdd %spec_3 %uint_3
%WorkGroupSize = OpSpecConstantComposite %v3uint %spec_3_op %spec_4 %uint_2
%main = OpFunction %void None %3
%5 = OpLabel
%20 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
%21 = OpLoad %uint %20
%24 = OpAccessChain %_ptr_Uniform_v4float %_ %int_0 %21
%25 = OpLoad %v4float %24
%26 = OpCompositeConstruct %v4float %float_2 %float_2 %float_2 %float_2
%27 = OpFAdd %v4float %25 %26
%wg_f = OpConvertUToF %v3float %WorkGroupSize
%wg_f4 = OpVectorShuffle %v4float %wg_f %wg_f 0 1 2 2
; Test that we can use the spec constants directly which needs to translate to gl_WorkGroupSize.elem.
; Needs special case handling.
%res = OpFAdd %v4float %27 %wg_f4
%f0 = OpConvertSToF %float %spec_3
%f1 = OpConvertSToF %float %spec_4
%f2 = OpConvertSToF %float %uint_2
%res1 = OpVectorTimesScalar %v4float %res %f0
%res2 = OpVectorTimesScalar %v4float %res1 %f1
%res3 = OpVectorTimesScalar %v4float %res2 %f2
%28 = OpAccessChain %_ptr_Uniform_v4float %_ %int_0 %21
OpStore %28 %res3
OpReturn
OpFunctionEnd

60
shaders-msl-no-opt/asm/comp/local-size-id-override.asm.comp Normal file
View File

@ -0,0 +1,60 @@
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID
OpExecutionModeId %main LocalSizeId %spec_3 %spec_4 %uint_2
OpSource GLSL 450
OpName %main "main"
OpName %SSBO "SSBO"
OpMemberName %SSBO 0 "values"
OpName %_ ""
OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
OpDecorate %_runtimearr_v4float ArrayStride 16
OpMemberDecorate %SSBO 0 Offset 0
OpDecorate %SSBO Block
OpDecorate %_ DescriptorSet 0
OpDecorate %_ Binding 0
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
OpDecorate %spec_1 SpecId 1
OpDecorate %spec_2 SpecId 2
OpDecorate %spec_3 SpecId 3
OpDecorate %spec_4 SpecId 4
OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
%void = OpTypeVoid
%3 = OpTypeFunction %void
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_runtimearr_v4float = OpTypeRuntimeArray %v4float
%SSBO = OpTypeStruct %_runtimearr_v4float
%_ptr_Uniform_SSBO = OpTypePointer StorageBuffer %SSBO
%_ = OpVariable %_ptr_Uniform_SSBO StorageBuffer
%int = OpTypeInt 32 1
%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
%float_2 = OpConstant %float 2
%_ptr_Uniform_v4float = OpTypePointer StorageBuffer %v4float
%spec_1 = OpSpecConstant %uint 11
%spec_2 = OpSpecConstant %uint 12
%spec_3 = OpSpecConstant %uint 13
%spec_4 = OpSpecConstant %uint 14
%uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%uint_3 = OpConstant %uint 3
%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %uint_3 %spec_1 %spec_2
%main = OpFunction %void None %3
%5 = OpLabel
%20 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
%21 = OpLoad %uint %20
%24 = OpAccessChain %_ptr_Uniform_v4float %_ %int_0 %21
%25 = OpLoad %v4float %24
%26 = OpCompositeConstruct %v4float %float_2 %float_2 %float_2 %float_2
%27 = OpFAdd %v4float %25 %26
%28 = OpAccessChain %_ptr_Uniform_v4float %_ %int_0 %21
OpStore %28 %27
OpReturn
OpFunctionEnd

76
shaders-msl-no-opt/asm/comp/local-size-id.asm.comp Normal file
View File

@ -0,0 +1,76 @@
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID
OpExecutionModeId %main LocalSizeId %spec_3 %spec_4 %uint_2
OpSource GLSL 450
OpName %main "main"
OpName %SSBO "SSBO"
OpMemberName %SSBO 0 "values"
OpName %_ ""
OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
OpDecorate %_runtimearr_v4float ArrayStride 16
OpMemberDecorate %SSBO 0 Offset 0
OpDecorate %SSBO Block
OpDecorate %_ DescriptorSet 0
OpDecorate %_ Binding 0
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
OpDecorate %spec_1 SpecId 1
OpDecorate %spec_2 SpecId 2
OpDecorate %spec_3 SpecId 3
OpDecorate %spec_4 SpecId 4
%void = OpTypeVoid
%3 = OpTypeFunction %void
%float = OpTypeFloat 32
%v3float = OpTypeVector %float 3
%v4float = OpTypeVector %float 4
%_runtimearr_v4float = OpTypeRuntimeArray %v4float
%SSBO = OpTypeStruct %_runtimearr_v4float
%_ptr_Uniform_SSBO = OpTypePointer StorageBuffer %SSBO
%_ = OpVariable %_ptr_Uniform_SSBO StorageBuffer
%int = OpTypeInt 32 1
%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
%float_2 = OpConstant %float 2
%_ptr_Uniform_v4float = OpTypePointer StorageBuffer %v4float
; Test that we can declare the spec constant as signed.
; Needs implicit bitcast since WorkGroupSize is uint.
%spec_1 = OpSpecConstant %int 11
%spec_2 = OpSpecConstant %int 12
%spec_3 = OpSpecConstant %int 13
%spec_4 = OpSpecConstant %int 14
%uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%uint_3 = OpConstant %uint 3
; Test that we can build spec constant composites out of local size id values.
; Needs special case handling.
%spec_3_op = OpSpecConstantOp %uint IAdd %spec_3 %uint_3
%WorkGroupSize = OpSpecConstantComposite %v3uint %spec_3_op %spec_4 %uint_2
%main = OpFunction %void None %3
%5 = OpLabel
%20 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
%21 = OpLoad %uint %20
%24 = OpAccessChain %_ptr_Uniform_v4float %_ %int_0 %21
%25 = OpLoad %v4float %24
%26 = OpCompositeConstruct %v4float %float_2 %float_2 %float_2 %float_2
%27 = OpFAdd %v4float %25 %26
%wg_f = OpConvertUToF %v3float %WorkGroupSize
%wg_f4 = OpVectorShuffle %v4float %wg_f %wg_f 0 1 2 2
; Test that we can use the spec constants directly which needs to translate to gl_WorkGroupSize.elem.
; Needs special case handling.
%res = OpFAdd %v4float %27 %wg_f4
%f0 = OpConvertSToF %float %spec_3
%f1 = OpConvertSToF %float %spec_4
%f2 = OpConvertSToF %float %uint_2
%res1 = OpVectorTimesScalar %v4float %res %f0
%res2 = OpVectorTimesScalar %v4float %res1 %f1
%res3 = OpVectorTimesScalar %v4float %res2 %f2
%28 = OpAccessChain %_ptr_Uniform_v4float %_ %int_0 %21
OpStore %28 %res3
OpReturn
OpFunctionEnd

60
shaders-no-opt/asm/comp/local-size-id-override.vk.asm.comp Normal file
View File

@ -0,0 +1,60 @@
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID
OpExecutionModeId %main LocalSizeId %spec_3 %spec_4 %uint_2
OpSource GLSL 450
OpName %main "main"
OpName %SSBO "SSBO"
OpMemberName %SSBO 0 "values"
OpName %_ ""
OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
OpDecorate %_runtimearr_v4float ArrayStride 16
OpMemberDecorate %SSBO 0 Offset 0
OpDecorate %SSBO Block
OpDecorate %_ DescriptorSet 0
OpDecorate %_ Binding 0
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
OpDecorate %spec_1 SpecId 1
OpDecorate %spec_2 SpecId 2
OpDecorate %spec_3 SpecId 3
OpDecorate %spec_4 SpecId 4
OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
%void = OpTypeVoid
%3 = OpTypeFunction %void
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_runtimearr_v4float = OpTypeRuntimeArray %v4float
%SSBO = OpTypeStruct %_runtimearr_v4float
%_ptr_Uniform_SSBO = OpTypePointer StorageBuffer %SSBO
%_ = OpVariable %_ptr_Uniform_SSBO StorageBuffer
%int = OpTypeInt 32 1
%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
%float_2 = OpConstant %float 2
%_ptr_Uniform_v4float = OpTypePointer StorageBuffer %v4float
%spec_1 = OpSpecConstant %uint 11
%spec_2 = OpSpecConstant %uint 12
%spec_3 = OpSpecConstant %uint 13
%spec_4 = OpSpecConstant %uint 14
%uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%uint_3 = OpConstant %uint 3
%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %uint_3 %spec_1 %spec_2
%main = OpFunction %void None %3
%5 = OpLabel
%20 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
%21 = OpLoad %uint %20
%24 = OpAccessChain %_ptr_Uniform_v4float %_ %int_0 %21
%25 = OpLoad %v4float %24
%26 = OpCompositeConstruct %v4float %float_2 %float_2 %float_2 %float_2
%27 = OpFAdd %v4float %25 %26
%28 = OpAccessChain %_ptr_Uniform_v4float %_ %int_0 %21
OpStore %28 %27
OpReturn
OpFunctionEnd

76
shaders-no-opt/asm/comp/local-size-id.vk.asm.comp Normal file
View File

@ -0,0 +1,76 @@
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID
OpExecutionModeId %main LocalSizeId %spec_3 %spec_4 %uint_2
OpSource GLSL 450
OpName %main "main"
OpName %SSBO "SSBO"
OpMemberName %SSBO 0 "values"
OpName %_ ""
OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
OpDecorate %_runtimearr_v4float ArrayStride 16
OpMemberDecorate %SSBO 0 Offset 0
OpDecorate %SSBO Block
OpDecorate %_ DescriptorSet 0
OpDecorate %_ Binding 0
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
OpDecorate %spec_1 SpecId 1
OpDecorate %spec_2 SpecId 2
OpDecorate %spec_3 SpecId 3
OpDecorate %spec_4 SpecId 4
%void = OpTypeVoid
%3 = OpTypeFunction %void
%float = OpTypeFloat 32
%v3float = OpTypeVector %float 3
%v4float = OpTypeVector %float 4
%_runtimearr_v4float = OpTypeRuntimeArray %v4float
%SSBO = OpTypeStruct %_runtimearr_v4float
%_ptr_Uniform_SSBO = OpTypePointer StorageBuffer %SSBO
%_ = OpVariable %_ptr_Uniform_SSBO StorageBuffer
%int = OpTypeInt 32 1
%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
%float_2 = OpConstant %float 2
%_ptr_Uniform_v4float = OpTypePointer StorageBuffer %v4float
; Test that we can declare the spec constant as signed.
; Needs implicit bitcast since WorkGroupSize is uint.
%spec_1 = OpSpecConstant %int 11
%spec_2 = OpSpecConstant %int 12
%spec_3 = OpSpecConstant %int 13
%spec_4 = OpSpecConstant %int 14
%uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%uint_3 = OpConstant %uint 3
; Test that we can build spec constant composites out of local size id values.
; Needs special case handling.
%spec_3_op = OpSpecConstantOp %uint IAdd %spec_3 %uint_3
%WorkGroupSize = OpSpecConstantComposite %v3uint %spec_3_op %spec_4 %uint_2
%main = OpFunction %void None %3
%5 = OpLabel
%20 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
%21 = OpLoad %uint %20
%24 = OpAccessChain %_ptr_Uniform_v4float %_ %int_0 %21
%25 = OpLoad %v4float %24
%26 = OpCompositeConstruct %v4float %float_2 %float_2 %float_2 %float_2
%27 = OpFAdd %v4float %25 %26
%wg_f = OpConvertUToF %v3float %WorkGroupSize
%wg_f4 = OpVectorShuffle %v4float %wg_f %wg_f 0 1 2 2
; Test that we can use the spec constants directly which needs to translate to gl_WorkGroupSize.elem.
; Needs special case handling.
%res = OpFAdd %v4float %27 %wg_f4
%f0 = OpConvertSToF %float %spec_3
%f1 = OpConvertSToF %float %spec_4
%f2 = OpConvertSToF %float %uint_2
%res1 = OpVectorTimesScalar %v4float %res %f0
%res2 = OpVectorTimesScalar %v4float %res1 %f1
%res3 = OpVectorTimesScalar %v4float %res2 %f2
%28 = OpAccessChain %_ptr_Uniform_v4float %_ %int_0 %21
OpStore %28 %res3
OpReturn
OpFunctionEnd

1
spirv_common.hpp
View File

@ -672,6 +672,7 @@ struct SPIREntryPoint
struct WorkgroupSize
{
uint32_t x = 0, y = 0, z = 0;
uint32_t id_x = 0, id_y = 0, id_z = 0;
uint32_t constant = 0; // Workgroup size can be expressed as a constant/spec-constant instead.
} workgroup_size;
uint32_t invocations = 0;

67
spirv_cross.cpp
View File

@ -1050,10 +1050,12 @@ void Compiler::parse_fixup()
if (id.get_type() == TypeConstant)
{
auto &c = id.get<SPIRConstant>();
if (ir.meta[c.self].decoration.builtin && ir.meta[c.self].decoration.builtin_type == BuiltInWorkgroupSize)
if (has_decoration(c.self, DecorationBuiltIn) &&
BuiltIn(get_decoration(c.self, DecorationBuiltIn)) == BuiltInWorkgroupSize)
{
// In current SPIR-V, there can be just one constant like this.
// All entry points will receive the constant value.
// WorkgroupSize take precedence over LocalSizeId.
for (auto &entry : ir.entry_points)
{
entry.second.workgroup_size.constant = c.self;
@ -2156,6 +2158,12 @@ void Compiler::set_execution_mode(ExecutionMode mode, uint32_t arg0, uint32_t ar
execution.workgroup_size.z = arg2;
break;
case ExecutionModeLocalSizeId:
execution.workgroup_size.id_x = arg0;
execution.workgroup_size.id_y = arg1;
execution.workgroup_size.id_z = arg2;
break;
case ExecutionModeInvocations:
execution.invocations = arg0;
break;
@ -2183,6 +2191,7 @@ uint32_t Compiler::get_work_group_size_specialization_constants(SpecializationCo
y = { 0, 0 };
z = { 0, 0 };
// WorkgroupSize builtin takes precedence over LocalSize / LocalSizeId.
if (execution.workgroup_size.constant != 0)
{
auto &c = get<SPIRConstant>(execution.workgroup_size.constant);
@ -2205,6 +2214,29 @@ uint32_t Compiler::get_work_group_size_specialization_constants(SpecializationCo
z.constant_id = get_decoration(c.m.c[0].id[2], DecorationSpecId);
}
}
else if (execution.flags.get(ExecutionModeLocalSizeId))
{
auto &cx = get<SPIRConstant>(execution.workgroup_size.id_x);
if (cx.specialization)
{
x.id = execution.workgroup_size.id_x;
x.constant_id = get_decoration(execution.workgroup_size.id_x, DecorationSpecId);
}
auto &cy = get<SPIRConstant>(execution.workgroup_size.id_y);
if (cy.specialization)
{
y.id = execution.workgroup_size.id_y;
y.constant_id = get_decoration(execution.workgroup_size.id_y, DecorationSpecId);
}
auto &cz = get<SPIRConstant>(execution.workgroup_size.id_z);
if (cz.specialization)
{
z.id = execution.workgroup_size.id_z;
z.constant_id = get_decoration(execution.workgroup_size.id_z, DecorationSpecId);
}
}
return execution.workgroup_size.constant;
}
@ -2214,15 +2246,42 @@ uint32_t Compiler::get_execution_mode_argument(spv::ExecutionMode mode, uint32_t
auto &execution = get_entry_point();
switch (mode)
{
case ExecutionModeLocalSizeId:
if (execution.flags.get(ExecutionModeLocalSizeId))
{
switch (index)
{
case 0:
return execution.workgroup_size.id_x;
case 1:
return execution.workgroup_size.id_y;
case 2:
return execution.workgroup_size.id_z;
default:
return 0;
}
}
else
return 0;
case ExecutionModeLocalSize:
switch (index)
{
case 0:
return execution.workgroup_size.x;
if (execution.flags.get(ExecutionModeLocalSizeId) && execution.workgroup_size.id_x != 0)
return get<SPIRConstant>(execution.workgroup_size.id_x).scalar();
else
return execution.workgroup_size.x;
case 1:
return execution.workgroup_size.y;
if (execution.flags.get(ExecutionModeLocalSizeId) && execution.workgroup_size.id_y != 0)
return get<SPIRConstant>(execution.workgroup_size.id_y).scalar();
else
return execution.workgroup_size.y;
case 2:
return execution.workgroup_size.z;
if (execution.flags.get(ExecutionModeLocalSizeId) && execution.workgroup_size.id_z != 0)
return get<SPIRConstant>(execution.workgroup_size.id_z).scalar();
else
return execution.workgroup_size.z;
default:
return 0;
}

7
spirv_cross.hpp
View File

@ -357,8 +357,11 @@ public:
void set_execution_mode(spv::ExecutionMode mode, uint32_t arg0 = 0, uint32_t arg1 = 0, uint32_t arg2 = 0);
// Gets argument for an execution mode (LocalSize, Invocations, OutputVertices).
// For LocalSize, the index argument is used to select the dimension (X = 0, Y = 1, Z = 2).
// For LocalSize or LocalSizeId, the index argument is used to select the dimension (X = 0, Y = 1, Z = 2).
// For execution modes which do not have arguments, 0 is returned.
// LocalSizeId query returns an ID. If LocalSizeId execution mode is not used, it returns 0.
// LocalSize always returns a literal. If execution mode is LocalSizeId,
// the literal (spec constant or not) is still returned.
uint32_t get_execution_mode_argument(spv::ExecutionMode mode, uint32_t index = 0) const;
spv::ExecutionModel get_execution_model() const;
@ -380,6 +383,8 @@ public:
// If the component is not a specialization constant, a zeroed out struct will be written.
// The return value is the constant ID of the builtin WorkGroupSize, but this is not expected to be useful
// for most use cases.
// If LocalSizeId is used, there is no uvec3 value representing the workgroup size, so the return value is 0,
// but x, y and z are written as normal if the components are specialization constants.
uint32_t get_work_group_size_specialization_constants(SpecializationConstant &x, SpecializationConstant &y,
SpecializationConstant &z) const;

76
spirv_glsl.cpp
View File

@ -637,6 +637,7 @@ string CompilerGLSL::compile()
backend.force_gl_in_out_block = true;
backend.supports_extensions = true;
backend.use_array_constructor = true;
backend.workgroup_size_is_hidden = true;
backend.support_precise_qualifier = (!options.es && options.version >= 400) || (options.es && options.version >= 320);
@ -707,6 +708,8 @@ void CompilerGLSL::build_workgroup_size(SmallVector<string> &arguments, const Sp
const SpecializationConstant &wg_y, const SpecializationConstant &wg_z)
{
auto &execution = get_entry_point();
bool builtin_workgroup = execution.workgroup_size.constant != 0;
bool use_local_size_id = !builtin_workgroup && execution.flags.get(ExecutionModeLocalSizeId);
if (wg_x.id)
{
@ -715,6 +718,8 @@ void CompilerGLSL::build_workgroup_size(SmallVector<string> &arguments, const Sp
else
arguments.push_back(join("local_size_x = ", get<SPIRConstant>(wg_x.id).specialization_constant_macro_name));
}
else if (use_local_size_id && execution.workgroup_size.id_x)
arguments.push_back(join("local_size_x = ", get<SPIRConstant>(execution.workgroup_size.id_x).scalar()));
else
arguments.push_back(join("local_size_x = ", execution.workgroup_size.x));
@ -725,6 +730,8 @@ void CompilerGLSL::build_workgroup_size(SmallVector<string> &arguments, const Sp
else
arguments.push_back(join("local_size_y = ", get<SPIRConstant>(wg_y.id).specialization_constant_macro_name));
}
else if (use_local_size_id && execution.workgroup_size.id_y)
arguments.push_back(join("local_size_y = ", get<SPIRConstant>(execution.workgroup_size.id_y).scalar()));
else
arguments.push_back(join("local_size_y = ", execution.workgroup_size.y));
@ -735,6 +742,8 @@ void CompilerGLSL::build_workgroup_size(SmallVector<string> &arguments, const Sp
else
arguments.push_back(join("local_size_z = ", get<SPIRConstant>(wg_z.id).specialization_constant_macro_name));
}
else if (use_local_size_id && execution.workgroup_size.id_z)
arguments.push_back(join("local_size_z = ", get<SPIRConstant>(execution.workgroup_size.id_z).scalar()));
else
arguments.push_back(join("local_size_z = ", execution.workgroup_size.z));
}
@ -1005,7 +1014,7 @@ void CompilerGLSL::emit_header()
case ExecutionModelGLCompute:
{
if (execution.workgroup_size.constant != 0)
if (execution.workgroup_size.constant != 0 || execution.flags.get(ExecutionModeLocalSizeId))
{
SpecializationConstant wg_x, wg_y, wg_z;
get_work_group_size_specialization_constants(wg_x, wg_y, wg_z);
@ -2673,6 +2682,26 @@ void CompilerGLSL::emit_specialization_constant_op(const SPIRConstantOp &constan
statement("const ", variable_decl(type, name), " = ", constant_op_expression(constant), ";");
}
int CompilerGLSL::get_constant_mapping_to_workgroup_component(const SPIRConstant &c) const
{
auto &entry_point = get_entry_point();
int index = -1;
// Need to redirect specialization constants which are used as WorkGroupSize to the builtin,
// since the spec constant declarations are never explicitly declared.
if (entry_point.workgroup_size.constant == 0 && entry_point.flags.get(ExecutionModeLocalSizeId))
{
if (c.self == entry_point.workgroup_size.id_x)
index = 0;
else if (c.self == entry_point.workgroup_size.id_y)
index = 1;
else if (c.self == entry_point.workgroup_size.id_z)
index = 2;
}
return index;
}
void CompilerGLSL::emit_constant(const SPIRConstant &constant)
{
auto &type = get<SPIRType>(constant.constant_type);
@ -3441,7 +3470,7 @@ void CompilerGLSL::emit_resources()
// If the work group size depends on a specialization constant, we need to declare the layout() block
// after constants (and their macros) have been declared.
if (execution.model == ExecutionModelGLCompute && !options.vulkan_semantics &&
execution.workgroup_size.constant != 0)
(execution.workgroup_size.constant != 0 || execution.flags.get(ExecutionModeLocalSizeId)))
{
SpecializationConstant wg_x, wg_y, wg_z;
get_work_group_size_specialization_constants(wg_x, wg_y, wg_z);
@ -4620,11 +4649,24 @@ string CompilerGLSL::to_expression(uint32_t id, bool register_expression_read)
auto &type = get<SPIRType>(c.constant_type);
// WorkGroupSize may be a constant.
auto &dec = ir.meta[c.self].decoration;
if (dec.builtin)
return builtin_to_glsl(dec.builtin_type, StorageClassGeneric);
if (has_decoration(c.self, DecorationBuiltIn))
return builtin_to_glsl(BuiltIn(get_decoration(c.self, DecorationBuiltIn)), StorageClassGeneric);
else if (c.specialization)
{
if (backend.workgroup_size_is_hidden)
{
int wg_index = get_constant_mapping_to_workgroup_component(c);
if (wg_index >= 0)
{
auto wg_size = join(builtin_to_glsl(BuiltInWorkgroupSize, StorageClassInput), vector_swizzle(1, wg_index));
if (type.basetype != SPIRType::UInt)
wg_size = bitcast_expression(type, SPIRType::UInt, wg_size);
return wg_size;
}
}
return to_name(id);
}
else if (c.is_used_as_lut)
return to_name(id);
else if (type.basetype == SPIRType::Struct && !backend.can_declare_struct_inline)
@ -5266,7 +5308,7 @@ string CompilerGLSL::constant_expression_vector(const SPIRConstant &c, uint32_t
for (uint32_t i = 0; i < c.vector_size(); i++)
{
if (c.vector_size() > 1 && c.specialization_constant_id(vector, i) != 0)
res += to_name(c.specialization_constant_id(vector, i));
res += to_expression(c.specialization_constant_id(vector, i));
else
res += convert_half_to_string(c, vector, i);
@ -5288,7 +5330,7 @@ string CompilerGLSL::constant_expression_vector(const SPIRConstant &c, uint32_t
for (uint32_t i = 0; i < c.vector_size(); i++)
{
if (c.vector_size() > 1 && c.specialization_constant_id(vector, i) != 0)
res += to_name(c.specialization_constant_id(vector, i));
res += to_expression(c.specialization_constant_id(vector, i));
else
res += convert_float_to_string(c, vector, i);
@ -5310,7 +5352,7 @@ string CompilerGLSL::constant_expression_vector(const SPIRConstant &c, uint32_t
for (uint32_t i = 0; i < c.vector_size(); i++)
{
if (c.vector_size() > 1 && c.specialization_constant_id(vector, i) != 0)
res += to_name(c.specialization_constant_id(vector, i));
res += to_expression(c.specialization_constant_id(vector, i));
else
res += convert_double_to_string(c, vector, i);
@ -5336,7 +5378,7 @@ string CompilerGLSL::constant_expression_vector(const SPIRConstant &c, uint32_t
for (uint32_t i = 0; i < c.vector_size(); i++)
{
if (c.vector_size() > 1 && c.specialization_constant_id(vector, i) != 0)
res += to_name(c.specialization_constant_id(vector, i));
res += to_expression(c.specialization_constant_id(vector, i));
else
res += convert_to_string(c.scalar_i64(vector, i), int64_type, backend.long_long_literal_suffix);
@ -5361,7 +5403,7 @@ string CompilerGLSL::constant_expression_vector(const SPIRConstant &c, uint32_t
for (uint32_t i = 0; i < c.vector_size(); i++)
{
if (c.vector_size() > 1 && c.specialization_constant_id(vector, i) != 0)
res += to_name(c.specialization_constant_id(vector, i));
res += to_expression(c.specialization_constant_id(vector, i));
else
{
res += convert_to_string(c.scalar_u64(vector, i));
@ -5396,7 +5438,7 @@ string CompilerGLSL::constant_expression_vector(const SPIRConstant &c, uint32_t
for (uint32_t i = 0; i < c.vector_size(); i++)
{
if (c.vector_size() > 1 && c.specialization_constant_id(vector, i) != 0)
res += to_name(c.specialization_constant_id(vector, i));
res += to_expression(c.specialization_constant_id(vector, i));
else
{
res += convert_to_string(c.scalar(vector, i));
@ -5426,7 +5468,7 @@ string CompilerGLSL::constant_expression_vector(const SPIRConstant &c, uint32_t
for (uint32_t i = 0; i < c.vector_size(); i++)
{
if (c.vector_size() > 1 && c.specialization_constant_id(vector, i) != 0)
res += to_name(c.specialization_constant_id(vector, i));
res += to_expression(c.specialization_constant_id(vector, i));
else
res += convert_to_string(c.scalar_i32(vector, i));
if (i + 1 < c.vector_size())
@ -5445,7 +5487,7 @@ string CompilerGLSL::constant_expression_vector(const SPIRConstant &c, uint32_t
for (uint32_t i = 0; i < c.vector_size(); i++)
{
if (c.vector_size() > 1 && c.specialization_constant_id(vector, i) != 0)
res += to_name(c.specialization_constant_id(vector, i));
res += to_expression(c.specialization_constant_id(vector, i));
else
{
if (*backend.uint16_t_literal_suffix)
@ -5479,7 +5521,7 @@ string CompilerGLSL::constant_expression_vector(const SPIRConstant &c, uint32_t
for (uint32_t i = 0; i < c.vector_size(); i++)
{
if (c.vector_size() > 1 && c.specialization_constant_id(vector, i) != 0)
res += to_name(c.specialization_constant_id(vector, i));
res += to_expression(c.specialization_constant_id(vector, i));
else
{
if (*backend.int16_t_literal_suffix)
@ -5513,7 +5555,7 @@ string CompilerGLSL::constant_expression_vector(const SPIRConstant &c, uint32_t
for (uint32_t i = 0; i < c.vector_size(); i++)
{
if (c.vector_size() > 1 && c.specialization_constant_id(vector, i) != 0)
res += to_name(c.specialization_constant_id(vector, i));
res += to_expression(c.specialization_constant_id(vector, i));
else
{
res += type_to_glsl(scalar_type);
@ -5538,7 +5580,7 @@ string CompilerGLSL::constant_expression_vector(const SPIRConstant &c, uint32_t
for (uint32_t i = 0; i < c.vector_size(); i++)
{
if (c.vector_size() > 1 && c.specialization_constant_id(vector, i) != 0)
res += to_name(c.specialization_constant_id(vector, i));
res += to_expression(c.specialization_constant_id(vector, i));
else
{
res += type_to_glsl(scalar_type);
@ -5561,7 +5603,7 @@ string CompilerGLSL::constant_expression_vector(const SPIRConstant &c, uint32_t
for (uint32_t i = 0; i < c.vector_size(); i++)
{
if (c.vector_size() > 1 && c.specialization_constant_id(vector, i) != 0)
res += to_name(c.specialization_constant_id(vector, i));
res += to_expression(c.specialization_constant_id(vector, i));
else
res += c.scalar(vector, i) ? "true" : "false";

2
spirv_glsl.hpp
View File

@ -587,6 +587,7 @@ protected:
bool support_pointer_to_pointer = false;
bool support_precise_qualifier = false;
bool support_64bit_switch = false;
bool workgroup_size_is_hidden = false;
} backend;
void emit_struct(SPIRType &type);
@ -610,6 +611,7 @@ protected:
void emit_block_chain(SPIRBlock &block);
void emit_hoisted_temporaries(SmallVector<std::pair<TypeID, ID>> &temporaries);
std::string constant_value_macro_name(uint32_t id);
int get_constant_mapping_to_workgroup_component(const SPIRConstant &constant) const;
void emit_constant(const SPIRConstant &constant);
void emit_specialization_constant_op(const SPIRConstantOp &constant);
std::string emit_continue_block(uint32_t continue_block, bool follow_true_block, bool follow_false_block);

30
spirv_hlsl.cpp
View File

@ -1215,14 +1215,20 @@ void CompilerHLSL::emit_specialization_constants_and_structs()
auto &type = get<SPIRType>(c.constant_type);
auto name = to_name(c.self);
// HLSL does not support specialization constants, so fallback to macros.
c.specialization_constant_macro_name =
constant_value_macro_name(get_decoration(c.self, DecorationSpecId));
if (has_decoration(c.self, DecorationSpecId))
{
// HLSL does not support specialization constants, so fallback to macros.
c.specialization_constant_macro_name =
constant_value_macro_name(get_decoration(c.self, DecorationSpecId));
statement("#ifndef ", c.specialization_constant_macro_name);
statement("#define ", c.specialization_constant_macro_name, " ", constant_expression(c));
statement("#endif");
statement("static const ", variable_decl(type, name), " = ", c.specialization_constant_macro_name, ";");
}
else
statement("static const ", variable_decl(type, name), " = ", constant_expression(c), ";");
statement("#ifndef ", c.specialization_constant_macro_name);
statement("#define ", c.specialization_constant_macro_name, " ", constant_expression(c));
statement("#endif");
statement("static const ", variable_decl(type, name), " = ", c.specialization_constant_macro_name, ";");
emitted = true;
}
}
@ -2431,6 +2437,16 @@ void CompilerHLSL::emit_hlsl_entry_point()
uint32_t y = execution.workgroup_size.y;
uint32_t z = execution.workgroup_size.z;
if (!execution.workgroup_size.constant && execution.flags.get(ExecutionModeLocalSizeId))
{
if (execution.workgroup_size.id_x)
x = get<SPIRConstant>(execution.workgroup_size.id_x).scalar();
if (execution.workgroup_size.id_y)
y = get<SPIRConstant>(execution.workgroup_size.id_y).scalar();
if (execution.workgroup_size.id_z)
z = get<SPIRConstant>(execution.workgroup_size.id_z).scalar();
}
auto x_expr = wg_x.id ? get<SPIRConstant>(wg_x.id).specialization_constant_macro_name : to_string(x);
auto y_expr = wg_y.id ? get<SPIRConstant>(wg_y.id).specialization_constant_macro_name : to_string(y);
auto z_expr = wg_z.id ? get<SPIRConstant>(wg_z.id).specialization_constant_macro_name : to_string(z);

16
spirv_parser.cpp
View File

@ -344,6 +344,22 @@ void Parser::parse(const Instruction &instruction)
break;
}
case OpExecutionModeId:
{
auto &execution = ir.entry_points[ops[0]];
auto mode = static_cast<ExecutionMode>(ops[1]);
execution.flags.set(mode);
if (mode == ExecutionModeLocalSizeId)
{
execution.workgroup_size.id_x = ops[2];
execution.workgroup_size.id_y = ops[3];
execution.workgroup_size.id_z = ops[4];
}
break;
}
case OpName:
{
uint32_t id = ops[0];