diff options
author | Hans-Kristian Arntzen <post@arntzen-software.no> | 2022-01-18 17:25:17 +0300 |
---|---|---|
committer | GitHub <noreply@github.com> | 2022-01-18 17:25:17 +0300 |
commit | 5a29181b357dc572125e7da2b8c24020be848af7 (patch) | |
tree | 2b2253b3486fd10b0126c9d4233a80114c1bdde9 /reference | |
parent | 08d5f5ed181f489e77f94d168071fb5605f2792a (diff) | |
parent | ac46140ba394ba5f85b6cfc72a5f5810084e8191 (diff) |
Merge pull request #1851 from KhronosGroup/fix-1835
Handle aliased names in spec constants.
Diffstat (limited to 'reference')
7 files changed, 260 insertions, 0 deletions
diff --git a/reference/shaders-hlsl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp b/reference/shaders-hlsl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp new file mode 100644 index 00000000..d3dc5337 --- /dev/null +++ b/reference/shaders-hlsl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp @@ -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(); +} diff --git a/reference/shaders-hlsl-no-opt/asm/comp/spec-constant-name-aliasing.asm.comp b/reference/shaders-hlsl-no-opt/asm/comp/spec-constant-name-aliasing.asm.comp new file mode 100644 index 00000000..423beee6 --- /dev/null +++ b/reference/shaders-hlsl-no-opt/asm/comp/spec-constant-name-aliasing.asm.comp @@ -0,0 +1,51 @@ +#ifndef SPIRV_CROSS_CONSTANT_ID_0 +#define SPIRV_CROSS_CONSTANT_ID_0 0 +#endif +static const int A = SPIRV_CROSS_CONSTANT_ID_0; +#ifndef SPIRV_CROSS_CONSTANT_ID_1 +#define SPIRV_CROSS_CONSTANT_ID_1 1 +#endif +static const int A_1 = SPIRV_CROSS_CONSTANT_ID_1; +#ifndef SPIRV_CROSS_CONSTANT_ID_2 +#define SPIRV_CROSS_CONSTANT_ID_2 2 +#endif +static const int A_2 = SPIRV_CROSS_CONSTANT_ID_2; +#ifndef SPIRV_CROSS_CONSTANT_ID_3 +#define SPIRV_CROSS_CONSTANT_ID_3 3 +#endif +static const int A_3 = SPIRV_CROSS_CONSTANT_ID_3; +#ifndef SPIRV_CROSS_CONSTANT_ID_4 +#define SPIRV_CROSS_CONSTANT_ID_4 4 +#endif +static const int A_4 = SPIRV_CROSS_CONSTANT_ID_4; +#ifndef SPIRV_CROSS_CONSTANT_ID_5 +#define SPIRV_CROSS_CONSTANT_ID_5 5 +#endif +static const int A_5 = SPIRV_CROSS_CONSTANT_ID_5; +static const int A_6 = (A - A_1); +static const int A_7 = (A_6 - A_2); +static const int A_8 = (A_7 - A_3); +static const int A_9 = (A_8 - A_4); +static const int A_10 = (A_9 - A_5); +static const int A_11 = (A_10 + A_5); +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + +RWByteAddressBuffer _5 : register(u0); + +static uint3 gl_GlobalInvocationID; +struct SPIRV_Cross_Input +{ + uint3 gl_GlobalInvocationID : SV_DispatchThreadID; +}; + +void comp_main() +{ + _5.Store(gl_GlobalInvocationID.x * 4 + 0, uint(A_11)); +} + +[numthreads(1, 1, 1)] +void main(SPIRV_Cross_Input stage_input) +{ + gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID; + comp_main(); +} diff --git a/reference/shaders-msl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp b/reference/shaders-msl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp new file mode 100644 index 00000000..a5b6fc32 --- /dev/null +++ b/reference/shaders-msl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp @@ -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]; +} + diff --git a/reference/shaders-msl-no-opt/asm/comp/spec-constant-name-aliasing.asm.comp b/reference/shaders-msl-no-opt/asm/comp/spec-constant-name-aliasing.asm.comp new file mode 100644 index 00000000..dda85050 --- /dev/null +++ b/reference/shaders-msl-no-opt/asm/comp/spec-constant-name-aliasing.asm.comp @@ -0,0 +1,35 @@ +#include <metal_stdlib> +#include <simd/simd.h> + +using namespace metal; + +struct SSBO +{ + int values[1]; +}; + +constant int A_tmp [[function_constant(0)]]; +constant int A = is_function_constant_defined(A_tmp) ? A_tmp : 0; +constant int A_1_tmp [[function_constant(1)]]; +constant int A_1 = is_function_constant_defined(A_1_tmp) ? A_1_tmp : 1; +constant int A_2_tmp [[function_constant(2)]]; +constant int A_2 = is_function_constant_defined(A_2_tmp) ? A_2_tmp : 2; +constant int A_3_tmp [[function_constant(3)]]; +constant int A_3 = is_function_constant_defined(A_3_tmp) ? A_3_tmp : 3; +constant int A_4_tmp [[function_constant(4)]]; +constant int A_4 = is_function_constant_defined(A_4_tmp) ? A_4_tmp : 4; +constant int A_5_tmp [[function_constant(5)]]; +constant int A_5 = is_function_constant_defined(A_5_tmp) ? A_5_tmp : 5; +constant int A_6 = (A - A_1); +constant int A_7 = (A_6 - A_2); +constant int A_8 = (A_7 - A_3); +constant int A_9 = (A_8 - A_4); +constant int A_10 = (A_9 - A_5); +constant int A_11 = (A_10 + A_5); +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + +kernel void main0(device SSBO& _5 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + _5.values[gl_GlobalInvocationID.x] = A_11; +} + diff --git a/reference/shaders-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp b/reference/shaders-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp new file mode 100644 index 00000000..1f43951a --- /dev/null +++ b/reference/shaders-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp @@ -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]; +} + diff --git a/reference/shaders-no-opt/asm/comp/spec-constant-name-aliasing.vk.asm.comp b/reference/shaders-no-opt/asm/comp/spec-constant-name-aliasing.vk.asm.comp new file mode 100644 index 00000000..e7b9dbf3 --- /dev/null +++ b/reference/shaders-no-opt/asm/comp/spec-constant-name-aliasing.vk.asm.comp @@ -0,0 +1,44 @@ +#version 450 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +#ifndef SPIRV_CROSS_CONSTANT_ID_0 +#define SPIRV_CROSS_CONSTANT_ID_0 0 +#endif +const int A = SPIRV_CROSS_CONSTANT_ID_0; +#ifndef SPIRV_CROSS_CONSTANT_ID_1 +#define SPIRV_CROSS_CONSTANT_ID_1 1 +#endif +const int A_1 = SPIRV_CROSS_CONSTANT_ID_1; +#ifndef SPIRV_CROSS_CONSTANT_ID_2 +#define SPIRV_CROSS_CONSTANT_ID_2 2 +#endif +const int A_2 = SPIRV_CROSS_CONSTANT_ID_2; +#ifndef SPIRV_CROSS_CONSTANT_ID_3 +#define SPIRV_CROSS_CONSTANT_ID_3 3 +#endif +const int A_3 = SPIRV_CROSS_CONSTANT_ID_3; +#ifndef SPIRV_CROSS_CONSTANT_ID_4 +#define SPIRV_CROSS_CONSTANT_ID_4 4 +#endif +const int A_4 = SPIRV_CROSS_CONSTANT_ID_4; +#ifndef SPIRV_CROSS_CONSTANT_ID_5 +#define SPIRV_CROSS_CONSTANT_ID_5 5 +#endif +const int A_5 = SPIRV_CROSS_CONSTANT_ID_5; +const int A_6 = (A - A_1); +const int A_7 = (A_6 - A_2); +const int A_8 = (A_7 - A_3); +const int A_9 = (A_8 - A_4); +const int A_10 = (A_9 - A_5); +const int A_11 = (A_10 + A_5); + +layout(binding = 0, std430) buffer SSBO +{ + int values[]; +} _5; + +void main() +{ + _5.values[gl_GlobalInvocationID.x] = A_11; +} + diff --git a/reference/shaders-no-opt/asm/comp/spec-constant-name-aliasing.vk.asm.comp.vk b/reference/shaders-no-opt/asm/comp/spec-constant-name-aliasing.vk.asm.comp.vk new file mode 100644 index 00000000..c31d0787 --- /dev/null +++ b/reference/shaders-no-opt/asm/comp/spec-constant-name-aliasing.vk.asm.comp.vk @@ -0,0 +1,26 @@ +#version 450 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(constant_id = 0) const int A = 0; +layout(constant_id = 1) const int A_1 = 1; +layout(constant_id = 2) const int A_2 = 2; +layout(constant_id = 3) const int A_3 = 3; +layout(constant_id = 4) const int A_4 = 4; +layout(constant_id = 5) const int A_5 = 5; +const int A_6 = (A - A_1); +const int A_7 = (A_6 - A_2); +const int A_8 = (A_7 - A_3); +const int A_9 = (A_8 - A_4); +const int A_10 = (A_9 - A_5); +const int A_11 = (A_10 + A_5); + +layout(set = 0, binding = 0, std430) buffer SSBO +{ + int values[]; +} _5; + +void main() +{ + _5.values[gl_GlobalInvocationID.x] = A_11; +} + |