EpicGames
/
SPIRV-Cross
зеркало из https://github.com/EpicGames/SPIRV-Cross.git
1
0
Форкнуть 0
Код Задачи Пакеты Проекты Релизы Вики Активность

Deal correctly with return sign of bitscan operations.

This commit is contained in:
Hans-Kristian Arntzen 2019-07-12 10:57:56 +02:00
Родитель 19ebbd48c7
Коммит 932ee0e328
13 изменённых файлов: 361 добавлений и 28 удалений
Показать статистику Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы

8
reference/opt/shaders-msl/vert/functions.vert
Просмотреть файл

@ -45,14 +45,14 @@ T degrees(T r)
// Implementation of the GLSL findLSB() function
template<typename T>
T findLSB(T x)
T spvFindLSB(T x)
{
return select(ctz(x), T(-1), x == T(0));
}
// Implementation of the signed GLSL findMSB() function
template<typename T>
T findSMSB(T x)
T spvFindSMSB(T x)
{
T v = select(x, T(-1) - x, x < T(0));
return select(clz(T(0)) - (clz(v) + T(1)), T(-1), v == T(0));
@ -112,8 +112,8 @@ vertex main0_out main0(main0_in in [[stage_in]], constant UBO& _18 [[buffer(0)]]
out.vNormal = in.aNormal;
out.vRotDeg = degrees(_18.rotRad);
out.vRotRad = radians(_18.rotDeg);
out.vLSB = findLSB(_18.bits);
out.vMSB = findSMSB(_18.bits);
out.vLSB = spvFindLSB(_18.bits);
out.vMSB = spvFindSMSB(_18.bits);
return out;
}

25
reference/shaders-hlsl-no-opt/asm/comp/bitscan.asm.comp Normal file
Просмотреть файл

@ -0,0 +1,25 @@
RWByteAddressBuffer _4 : register(u0);
void comp_main()
{
uint4 _19 = _4.Load4(0);
int4 _20 = int4(_4.Load4(16));
_4.Store4(0, firstbitlow(_19));
_4.Store4(16, uint4(int4(firstbitlow(_19))));
_4.Store4(0, uint4(firstbitlow(_20)));
_4.Store4(16, uint4(firstbitlow(_20)));
_4.Store4(0, firstbithigh(_19));
_4.Store4(16, uint4(int4(firstbithigh(_19))));
_4.Store4(0, firstbithigh(uint4(_20)));
_4.Store4(16, uint4(int4(firstbithigh(uint4(_20)))));
_4.Store4(0, uint4(firstbithigh(int4(_19))));
_4.Store4(16, uint4(firstbithigh(int4(_19))));
_4.Store4(0, uint4(firstbithigh(_20)));
_4.Store4(16, uint4(firstbithigh(_20)));
}
[numthreads(1, 1, 1)]
void main()
{
comp_main();
}

53
reference/shaders-msl-no-opt/asm/comp/bitscan.asm.comp Normal file
Просмотреть файл

@ -0,0 +1,53 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct SSBO
{
uint4 u;
int4 i;
};
// Implementation of the GLSL findLSB() function
template<typename T>
T spvFindLSB(T x)
{
return select(ctz(x), T(-1), x == T(0));
}
// Implementation of the signed GLSL findMSB() function
template<typename T>
T spvFindSMSB(T x)
{
T v = select(x, T(-1) - x, x < T(0));
return select(clz(T(0)) - (clz(v) + T(1)), T(-1), v == T(0));
}
// Implementation of the unsigned GLSL findMSB() function
template<typename T>
T spvFindUMSB(T x)
{
return select(clz(T(0)) - (clz(x) + T(1)), T(-1), x == T(0));
}
kernel void main0(device SSBO& _4 [[buffer(0)]])
{
uint4 _19 = _4.u;
int4 _20 = _4.i;
_4.u = spvFindLSB(_19);
_4.i = int4(spvFindLSB(_19));
_4.u = uint4(spvFindLSB(_20));
_4.i = spvFindLSB(_20);
_4.u = spvFindUMSB(_19);
_4.i = int4(spvFindUMSB(_19));
_4.u = spvFindUMSB(uint4(_20));
_4.i = int4(spvFindUMSB(uint4(_20)));
_4.u = uint4(spvFindSMSB(int4(_19)));
_4.i = spvFindSMSB(int4(_19));
_4.u = uint4(spvFindSMSB(_20));
_4.i = spvFindSMSB(_20);
}

12
reference/shaders-msl-no-opt/asm/comp/glsl-signed-operations.asm.comp
Просмотреть файл

@ -13,7 +13,7 @@ struct SSBO
// Implementation of the signed GLSL findMSB() function
template<typename T>
T findSMSB(T x)
T spvFindSMSB(T x)
{
T v = select(x, T(-1) - x, x < T(0));
return select(clz(T(0)) - (clz(v) + T(1)), T(-1), v == T(0));
@ -21,7 +21,7 @@ T findSMSB(T x)
// Implementation of the unsigned GLSL findMSB() function
template<typename T>
T findUMSB(T x)
T spvFindUMSB(T x)
{
return select(clz(T(0)) - (clz(x) + T(1)), T(-1), x == T(0));
}
@ -45,10 +45,10 @@ kernel void main0(device SSBO& _4 [[buffer(0)]])
_4.uints = uint4(sign(_19));
_4.ints = sign(int4(_20));
_4.uints = uint4(sign(int4(_20)));
_4.ints = findSMSB(int4(_20));
_4.uints = uint4(findSMSB(int4(_20)));
_4.ints = int4(findUMSB(uint4(_19)));
_4.uints = findUMSB(uint4(_19));
_4.ints = spvFindSMSB(int4(_20));
_4.uints = uint4(spvFindSMSB(int4(_20)));
_4.ints = int4(spvFindUMSB(uint4(_19)));
_4.uints = spvFindUMSB(uint4(_19));
_4.ints = min(_19, _19);
_4.uints = uint4(min(_19, int4(_20)));
_4.ints = min(int4(_20), int4(_20));

14
reference/shaders-msl-no-opt/comp/bitfield.comp
Просмотреть файл

@ -7,14 +7,14 @@ using namespace metal;
// Implementation of the GLSL findLSB() function
template<typename T>
T findLSB(T x)
T spvFindLSB(T x)
{
return select(ctz(x), T(-1), x == T(0));
}
// Implementation of the signed GLSL findMSB() function
template<typename T>
T findSMSB(T x)
T spvFindSMSB(T x)
{
T v = select(x, T(-1) - x, x < T(0));
return select(clz(T(0)) - (clz(v) + T(1)), T(-1), v == T(0));
@ -22,7 +22,7 @@ T findSMSB(T x)
// Implementation of the unsigned GLSL findMSB() function
template<typename T>
T findUMSB(T x)
T spvFindUMSB(T x)
{
return select(clz(T(0)) - (clz(x) + T(1)), T(-1), x == T(0));
}
@ -39,9 +39,9 @@ kernel void main0()
s = reverse_bits(s);
int v0 = popcount(u);
int v1 = popcount(s);
int v2 = int(findUMSB(u));
int v3 = findSMSB(s);
int v4 = findLSB(u);
int v5 = findLSB(s);
int v2 = int(spvFindUMSB(u));
int v3 = spvFindSMSB(s);
int v4 = int(spvFindLSB(u));
int v5 = spvFindLSB(s);
}

8
reference/shaders-msl/vert/functions.vert
Просмотреть файл

@ -45,14 +45,14 @@ T degrees(T r)
// Implementation of the GLSL findLSB() function
template<typename T>
T findLSB(T x)
T spvFindLSB(T x)
{
return select(ctz(x), T(-1), x == T(0));
}
// Implementation of the signed GLSL findMSB() function
template<typename T>
T findSMSB(T x)
T spvFindSMSB(T x)
{
T v = select(x, T(-1) - x, x < T(0));
return select(clz(T(0)) - (clz(v) + T(1)), T(-1), v == T(0));
@ -112,8 +112,8 @@ vertex main0_out main0(main0_in in [[stage_in]], constant UBO& _18 [[buffer(0)]]
out.vNormal = in.aNormal;
out.vRotDeg = degrees(_18.rotRad);
out.vRotRad = radians(_18.rotDeg);
out.vLSB = findLSB(_18.bits);
out.vMSB = findSMSB(_18.bits);
out.vLSB = spvFindLSB(_18.bits);
out.vMSB = spvFindSMSB(_18.bits);
return out;
}

27
reference/shaders-no-opt/asm/comp/bitscan.asm.comp Normal file
Просмотреть файл

@ -0,0 +1,27 @@
#version 450
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(binding = 0, std430) buffer SSBO
{
uvec4 u;
ivec4 i;
} _4;
void main()
{
uvec4 _19 = _4.u;
ivec4 _20 = _4.i;
_4.u = uvec4(findLSB(_19));
_4.i = findLSB(_19);
_4.u = uvec4(findLSB(_20));
_4.i = findLSB(_20);
_4.u = uvec4(findMSB(_19));
_4.i = findMSB(_19);
_4.u = uvec4(findMSB(uvec4(_20)));
_4.i = findMSB(uvec4(_20));
_4.u = uvec4(findMSB(ivec4(_19)));
_4.i = findMSB(ivec4(_19));
_4.u = uvec4(findMSB(_20));
_4.i = findMSB(_20);
}

72
shaders-hlsl-no-opt/asm/comp/bitscan.asm.comp Normal file
Просмотреть файл

@ -0,0 +1,72 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 7
; Bound: 35
; 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 %SSBO "SSBO"
OpMemberName %SSBO 0 "u"
OpMemberName %SSBO 1 "i"
OpMemberDecorate %SSBO 0 Offset 0
OpMemberDecorate %SSBO 1 Offset 16
OpDecorate %SSBO BufferBlock
OpDecorate %_ DescriptorSet 0
OpDecorate %_ Binding 0
%void = OpTypeVoid
%3 = OpTypeFunction %void
%int = OpTypeInt 32 1
%ivec4 = OpTypeVector %int 4
%uint = OpTypeInt 32 0
%uvec4 = OpTypeVector %uint 4
%SSBO = OpTypeStruct %uvec4 %ivec4
%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
%_ = OpVariable %_ptr_Uniform_SSBO Uniform
%int_0 = OpConstant %int 0
%_ptr_Uniform_uvec4 = OpTypePointer Uniform %uvec4
%int_1 = OpConstant %int 1
%_ptr_Uniform_ivec4 = OpTypePointer Uniform %ivec4
%main = OpFunction %void None %3
%5 = OpLabel
%uptr = OpAccessChain %_ptr_Uniform_uvec4 %_ %int_0
%iptr = OpAccessChain %_ptr_Uniform_ivec4 %_ %int_1
%uvalue = OpLoad %uvec4 %uptr
%ivalue = OpLoad %ivec4 %iptr
%lsb_uint_to_uint = OpExtInst %uvec4 %1 FindILsb %uvalue
%lsb_uint_to_int = OpExtInst %ivec4 %1 FindILsb %uvalue
%lsb_int_to_uint = OpExtInst %uvec4 %1 FindILsb %ivalue
%lsb_int_to_int = OpExtInst %ivec4 %1 FindILsb %ivalue
%umsb_uint_to_uint = OpExtInst %uvec4 %1 FindUMsb %uvalue
%umsb_uint_to_int = OpExtInst %ivec4 %1 FindUMsb %uvalue
%umsb_int_to_uint = OpExtInst %uvec4 %1 FindUMsb %ivalue
%umsb_int_to_int = OpExtInst %ivec4 %1 FindUMsb %ivalue
%smsb_uint_to_uint = OpExtInst %uvec4 %1 FindSMsb %uvalue
%smsb_uint_to_int = OpExtInst %ivec4 %1 FindSMsb %uvalue
%smsb_int_to_uint = OpExtInst %uvec4 %1 FindSMsb %ivalue
%smsb_int_to_int = OpExtInst %ivec4 %1 FindSMsb %ivalue
OpStore %uptr %lsb_uint_to_uint
OpStore %iptr %lsb_uint_to_int
OpStore %uptr %lsb_int_to_uint
OpStore %iptr %lsb_int_to_int
OpStore %uptr %umsb_uint_to_uint
OpStore %iptr %umsb_uint_to_int
OpStore %uptr %umsb_int_to_uint
OpStore %iptr %umsb_int_to_int
OpStore %uptr %smsb_uint_to_uint
OpStore %iptr %smsb_uint_to_int
OpStore %uptr %smsb_int_to_uint
OpStore %iptr %smsb_int_to_int
OpReturn
OpFunctionEnd

72
shaders-msl-no-opt/asm/comp/bitscan.asm.comp Normal file
Просмотреть файл

@ -0,0 +1,72 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 7
; Bound: 35
; 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 %SSBO "SSBO"
OpMemberName %SSBO 0 "u"
OpMemberName %SSBO 1 "i"
OpMemberDecorate %SSBO 0 Offset 0
OpMemberDecorate %SSBO 1 Offset 16
OpDecorate %SSBO BufferBlock
OpDecorate %_ DescriptorSet 0
OpDecorate %_ Binding 0
%void = OpTypeVoid
%3 = OpTypeFunction %void
%int = OpTypeInt 32 1
%ivec4 = OpTypeVector %int 4
%uint = OpTypeInt 32 0
%uvec4 = OpTypeVector %uint 4
%SSBO = OpTypeStruct %uvec4 %ivec4
%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
%_ = OpVariable %_ptr_Uniform_SSBO Uniform
%int_0 = OpConstant %int 0
%_ptr_Uniform_uvec4 = OpTypePointer Uniform %uvec4
%int_1 = OpConstant %int 1
%_ptr_Uniform_ivec4 = OpTypePointer Uniform %ivec4
%main = OpFunction %void None %3
%5 = OpLabel
%uptr = OpAccessChain %_ptr_Uniform_uvec4 %_ %int_0
%iptr = OpAccessChain %_ptr_Uniform_ivec4 %_ %int_1
%uvalue = OpLoad %uvec4 %uptr
%ivalue = OpLoad %ivec4 %iptr
%lsb_uint_to_uint = OpExtInst %uvec4 %1 FindILsb %uvalue
%lsb_uint_to_int = OpExtInst %ivec4 %1 FindILsb %uvalue
%lsb_int_to_uint = OpExtInst %uvec4 %1 FindILsb %ivalue
%lsb_int_to_int = OpExtInst %ivec4 %1 FindILsb %ivalue
%umsb_uint_to_uint = OpExtInst %uvec4 %1 FindUMsb %uvalue
%umsb_uint_to_int = OpExtInst %ivec4 %1 FindUMsb %uvalue
%umsb_int_to_uint = OpExtInst %uvec4 %1 FindUMsb %ivalue
%umsb_int_to_int = OpExtInst %ivec4 %1 FindUMsb %ivalue
%smsb_uint_to_uint = OpExtInst %uvec4 %1 FindSMsb %uvalue
%smsb_uint_to_int = OpExtInst %ivec4 %1 FindSMsb %uvalue
%smsb_int_to_uint = OpExtInst %uvec4 %1 FindSMsb %ivalue
%smsb_int_to_int = OpExtInst %ivec4 %1 FindSMsb %ivalue
OpStore %uptr %lsb_uint_to_uint
OpStore %iptr %lsb_uint_to_int
OpStore %uptr %lsb_int_to_uint
OpStore %iptr %lsb_int_to_int
OpStore %uptr %umsb_uint_to_uint
OpStore %iptr %umsb_uint_to_int
OpStore %uptr %umsb_int_to_uint
OpStore %iptr %umsb_int_to_int
OpStore %uptr %smsb_uint_to_uint
OpStore %iptr %smsb_uint_to_int
OpStore %uptr %smsb_int_to_uint
OpStore %iptr %smsb_int_to_int
OpReturn
OpFunctionEnd

72
shaders-no-opt/asm/comp/bitscan.asm.comp Normal file
Просмотреть файл

@ -0,0 +1,72 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 7
; Bound: 35
; 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 %SSBO "SSBO"
OpMemberName %SSBO 0 "u"
OpMemberName %SSBO 1 "i"
OpMemberDecorate %SSBO 0 Offset 0
OpMemberDecorate %SSBO 1 Offset 16
OpDecorate %SSBO BufferBlock
OpDecorate %_ DescriptorSet 0
OpDecorate %_ Binding 0
%void = OpTypeVoid
%3 = OpTypeFunction %void
%int = OpTypeInt 32 1
%ivec4 = OpTypeVector %int 4
%uint = OpTypeInt 32 0
%uvec4 = OpTypeVector %uint 4
%SSBO = OpTypeStruct %uvec4 %ivec4
%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
%_ = OpVariable %_ptr_Uniform_SSBO Uniform
%int_0 = OpConstant %int 0
%_ptr_Uniform_uvec4 = OpTypePointer Uniform %uvec4
%int_1 = OpConstant %int 1
%_ptr_Uniform_ivec4 = OpTypePointer Uniform %ivec4
%main = OpFunction %void None %3
%5 = OpLabel
%uptr = OpAccessChain %_ptr_Uniform_uvec4 %_ %int_0
%iptr = OpAccessChain %_ptr_Uniform_ivec4 %_ %int_1
%uvalue = OpLoad %uvec4 %uptr
%ivalue = OpLoad %ivec4 %iptr
%lsb_uint_to_uint = OpExtInst %uvec4 %1 FindILsb %uvalue
%lsb_uint_to_int = OpExtInst %ivec4 %1 FindILsb %uvalue
%lsb_int_to_uint = OpExtInst %uvec4 %1 FindILsb %ivalue
%lsb_int_to_int = OpExtInst %ivec4 %1 FindILsb %ivalue
%umsb_uint_to_uint = OpExtInst %uvec4 %1 FindUMsb %uvalue
%umsb_uint_to_int = OpExtInst %ivec4 %1 FindUMsb %uvalue
%umsb_int_to_uint = OpExtInst %uvec4 %1 FindUMsb %ivalue
%umsb_int_to_int = OpExtInst %ivec4 %1 FindUMsb %ivalue
%smsb_uint_to_uint = OpExtInst %uvec4 %1 FindSMsb %uvalue
%smsb_uint_to_int = OpExtInst %ivec4 %1 FindSMsb %uvalue
%smsb_int_to_uint = OpExtInst %uvec4 %1 FindSMsb %ivalue
%smsb_int_to_int = OpExtInst %ivec4 %1 FindSMsb %ivalue
OpStore %uptr %lsb_uint_to_uint
OpStore %iptr %lsb_uint_to_int
OpStore %uptr %lsb_int_to_uint
OpStore %iptr %lsb_int_to_int
OpStore %uptr %umsb_uint_to_uint
OpStore %iptr %umsb_uint_to_int
OpStore %uptr %umsb_int_to_uint
OpStore %iptr %umsb_int_to_int
OpStore %uptr %smsb_uint_to_uint
OpStore %iptr %smsb_uint_to_int
OpStore %uptr %smsb_int_to_uint
OpStore %iptr %smsb_int_to_int
OpReturn
OpFunctionEnd

3
spirv_glsl.cpp
Просмотреть файл

@ -5445,7 +5445,8 @@ void CompilerGLSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop,
// Bit-fiddling
case GLSLstd450FindILsb:
emit_unary_func_op(result_type, id, args[0], "findLSB");
// findLSB always returns int.
emit_unary_func_op_cast(result_type, id, args[0], "findLSB", expression_type(args[0]).basetype, int_type);
break;
case GLSLstd450FindSMsb:

5
spirv_hlsl.cpp
Просмотреть файл

@ -3237,8 +3237,11 @@ void CompilerHLSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop,
SPIRV_CROSS_THROW("packDouble2x32/unpackDouble2x32 not supported in HLSL.");
case GLSLstd450FindILsb:
emit_unary_func_op(result_type, id, args[0], "firstbitlow");
{
auto basetype = expression_type(args[0]).basetype;
emit_unary_func_op_cast(result_type, id, args[0], "firstbitlow", basetype, basetype);
break;
}
case GLSLstd450FindSMsb:
emit_unary_func_op_cast(result_type, id, args[0], "firstbithigh", int_type, int_type);

18
spirv_msl.cpp
Просмотреть файл

@ -2787,7 +2787,7 @@ void CompilerMSL::emit_custom_functions()
case SPVFuncImplFindILsb:
statement("// Implementation of the GLSL findLSB() function");
statement("template<typename T>");
statement("T findLSB(T x)");
statement("T spvFindLSB(T x)");
begin_scope();
statement("return select(ctz(x), T(-1), x == T(0));");
end_scope();
@ -2797,7 +2797,7 @@ void CompilerMSL::emit_custom_functions()
case SPVFuncImplFindUMsb:
statement("// Implementation of the unsigned GLSL findMSB() function");
statement("template<typename T>");
statement("T findUMSB(T x)");
statement("T spvFindUMSB(T x)");
begin_scope();
statement("return select(clz(T(0)) - (clz(x) + T(1)), T(-1), x == T(0));");
end_scope();
@ -2807,7 +2807,7 @@ void CompilerMSL::emit_custom_functions()
case SPVFuncImplFindSMsb:
statement("// Implementation of the signed GLSL findMSB() function");
statement("template<typename T>");
statement("T findSMSB(T x)");
statement("T spvFindSMSB(T x)");
begin_scope();
statement("T v = select(x, T(-1) - x, x < T(0));");
statement("return select(clz(T(0)) - (clz(v) + T(1)), T(-1), v == T(0));");
@ -4688,12 +4688,20 @@ void CompilerMSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop,
emit_unary_func_op(result_type, id, args[0], "rint");
break;
case GLSLstd450FindILsb:
{
// In this template version of findLSB, we return T.
auto basetype = expression_type(args[0]).basetype;
emit_unary_func_op_cast(result_type, id, args[0], "spvFindLSB", basetype, basetype);
break;
}
case GLSLstd450FindSMsb:
emit_unary_func_op_cast(result_type, id, args[0], "findSMSB", int_type, int_type);
emit_unary_func_op_cast(result_type, id, args[0], "spvFindSMSB", int_type, int_type);
break;
case GLSLstd450FindUMsb:
emit_unary_func_op_cast(result_type, id, args[0], "findUMSB", uint_type, uint_type);
emit_unary_func_op_cast(result_type, id, args[0], "spvFindUMSB", uint_type, uint_type);
break;
case GLSLstd450PackSnorm4x8: