HLSL: Fix unrolled S/G LE/LT/GE/GT opcodes.

Need to bitcast the unrolled expressions as well.
This commit is contained in:
Hans-Kristian Arntzen 2019-10-14 16:08:39 +02:00
parent 3bf9fa7ed6
commit a9be92569f
12 changed files with 523 additions and 26 deletions

View File

@ -0,0 +1,28 @@
RWByteAddressBuffer _5 : register(u0);
RWByteAddressBuffer _6 : register(u1);
void comp_main()
{
bool4 _31 = bool4(int(_5.Load4(16).x) < int4(_5.Load4(0)).x, int(_5.Load4(16).y) < int4(_5.Load4(0)).y, int(_5.Load4(16).z) < int4(_5.Load4(0)).z, int(_5.Load4(16).w) < int4(_5.Load4(0)).w);
bool4 _32 = bool4(int(_5.Load4(16).x) <= int4(_5.Load4(0)).x, int(_5.Load4(16).y) <= int4(_5.Load4(0)).y, int(_5.Load4(16).z) <= int4(_5.Load4(0)).z, int(_5.Load4(16).w) <= int4(_5.Load4(0)).w);
bool4 _33 = bool4(_5.Load4(16).x < uint(int4(_5.Load4(0)).x), _5.Load4(16).y < uint(int4(_5.Load4(0)).y), _5.Load4(16).z < uint(int4(_5.Load4(0)).z), _5.Load4(16).w < uint(int4(_5.Load4(0)).w));
bool4 _34 = bool4(_5.Load4(16).x <= uint(int4(_5.Load4(0)).x), _5.Load4(16).y <= uint(int4(_5.Load4(0)).y), _5.Load4(16).z <= uint(int4(_5.Load4(0)).z), _5.Load4(16).w <= uint(int4(_5.Load4(0)).w));
bool4 _35 = bool4(int(_5.Load4(16).x) > int4(_5.Load4(0)).x, int(_5.Load4(16).y) > int4(_5.Load4(0)).y, int(_5.Load4(16).z) > int4(_5.Load4(0)).z, int(_5.Load4(16).w) > int4(_5.Load4(0)).w);
bool4 _36 = bool4(int(_5.Load4(16).x) >= int4(_5.Load4(0)).x, int(_5.Load4(16).y) >= int4(_5.Load4(0)).y, int(_5.Load4(16).z) >= int4(_5.Load4(0)).z, int(_5.Load4(16).w) >= int4(_5.Load4(0)).w);
bool4 _37 = bool4(_5.Load4(16).x > uint(int4(_5.Load4(0)).x), _5.Load4(16).y > uint(int4(_5.Load4(0)).y), _5.Load4(16).z > uint(int4(_5.Load4(0)).z), _5.Load4(16).w > uint(int4(_5.Load4(0)).w));
bool4 _38 = bool4(_5.Load4(16).x >= uint(int4(_5.Load4(0)).x), _5.Load4(16).y >= uint(int4(_5.Load4(0)).y), _5.Load4(16).z >= uint(int4(_5.Load4(0)).z), _5.Load4(16).w >= uint(int4(_5.Load4(0)).w));
_6.Store4(0, uint4(_31.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _31.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _31.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _31.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
_6.Store4(0, uint4(_32.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _32.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _32.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _32.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
_6.Store4(0, uint4(_33.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _33.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _33.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _33.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
_6.Store4(0, uint4(_34.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _34.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _34.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _34.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
_6.Store4(0, uint4(_35.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _35.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _35.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _35.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
_6.Store4(0, uint4(_36.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _36.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _36.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _36.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
_6.Store4(0, uint4(_37.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _37.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _37.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _37.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
_6.Store4(0, uint4(_38.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _38.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _38.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _38.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
}
[numthreads(1, 1, 1)]
void main()
{
comp_main();
}

View File

@ -0,0 +1,29 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct _3
{
int4 _m0;
uint4 _m1;
};
struct _4
{
uint4 _m0;
int4 _m1;
};
kernel void main0(device _3& restrict _5 [[buffer(0)]], device _4& restrict _6 [[buffer(1)]])
{
_6._m0 = select(uint4(0u), uint4(1u), int4(_5._m1) < _5._m0);
_6._m0 = select(uint4(0u), uint4(1u), int4(_5._m1) <= _5._m0);
_6._m0 = select(uint4(0u), uint4(1u), _5._m1 < uint4(_5._m0));
_6._m0 = select(uint4(0u), uint4(1u), _5._m1 <= uint4(_5._m0));
_6._m0 = select(uint4(0u), uint4(1u), int4(_5._m1) > _5._m0);
_6._m0 = select(uint4(0u), uint4(1u), int4(_5._m1) >= _5._m0);
_6._m0 = select(uint4(0u), uint4(1u), _5._m1 > uint4(_5._m0));
_6._m0 = select(uint4(0u), uint4(1u), _5._m1 >= uint4(_5._m0));
}

View File

@ -0,0 +1,27 @@
#version 310 es
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(binding = 0, std430) restrict buffer _3_5
{
ivec4 _m0;
uvec4 _m1;
} _5;
layout(binding = 1, std430) restrict buffer _4_6
{
uvec4 _m0;
ivec4 _m1;
} _6;
void main()
{
_6._m0 = mix(uvec4(0u), uvec4(1u), lessThan(ivec4(_5._m1), _5._m0));
_6._m0 = mix(uvec4(0u), uvec4(1u), lessThanEqual(ivec4(_5._m1), _5._m0));
_6._m0 = mix(uvec4(0u), uvec4(1u), lessThan(_5._m1, uvec4(_5._m0)));
_6._m0 = mix(uvec4(0u), uvec4(1u), lessThanEqual(_5._m1, uvec4(_5._m0)));
_6._m0 = mix(uvec4(0u), uvec4(1u), greaterThan(ivec4(_5._m1), _5._m0));
_6._m0 = mix(uvec4(0u), uvec4(1u), greaterThanEqual(ivec4(_5._m1), _5._m0));
_6._m0 = mix(uvec4(0u), uvec4(1u), greaterThan(_5._m1, uvec4(_5._m0)));
_6._m0 = mix(uvec4(0u), uvec4(1u), greaterThanEqual(_5._m1, uvec4(_5._m0)));
}

View File

@ -0,0 +1,28 @@
RWByteAddressBuffer _5 : register(u0);
RWByteAddressBuffer _6 : register(u1);
void comp_main()
{
bool4 _31 = bool4(int(_5.Load4(16).x) < int4(_5.Load4(0)).x, int(_5.Load4(16).y) < int4(_5.Load4(0)).y, int(_5.Load4(16).z) < int4(_5.Load4(0)).z, int(_5.Load4(16).w) < int4(_5.Load4(0)).w);
bool4 _32 = bool4(int(_5.Load4(16).x) <= int4(_5.Load4(0)).x, int(_5.Load4(16).y) <= int4(_5.Load4(0)).y, int(_5.Load4(16).z) <= int4(_5.Load4(0)).z, int(_5.Load4(16).w) <= int4(_5.Load4(0)).w);
bool4 _33 = bool4(_5.Load4(16).x < uint(int4(_5.Load4(0)).x), _5.Load4(16).y < uint(int4(_5.Load4(0)).y), _5.Load4(16).z < uint(int4(_5.Load4(0)).z), _5.Load4(16).w < uint(int4(_5.Load4(0)).w));
bool4 _34 = bool4(_5.Load4(16).x <= uint(int4(_5.Load4(0)).x), _5.Load4(16).y <= uint(int4(_5.Load4(0)).y), _5.Load4(16).z <= uint(int4(_5.Load4(0)).z), _5.Load4(16).w <= uint(int4(_5.Load4(0)).w));
bool4 _35 = bool4(int(_5.Load4(16).x) > int4(_5.Load4(0)).x, int(_5.Load4(16).y) > int4(_5.Load4(0)).y, int(_5.Load4(16).z) > int4(_5.Load4(0)).z, int(_5.Load4(16).w) > int4(_5.Load4(0)).w);
bool4 _36 = bool4(int(_5.Load4(16).x) >= int4(_5.Load4(0)).x, int(_5.Load4(16).y) >= int4(_5.Load4(0)).y, int(_5.Load4(16).z) >= int4(_5.Load4(0)).z, int(_5.Load4(16).w) >= int4(_5.Load4(0)).w);
bool4 _37 = bool4(_5.Load4(16).x > uint(int4(_5.Load4(0)).x), _5.Load4(16).y > uint(int4(_5.Load4(0)).y), _5.Load4(16).z > uint(int4(_5.Load4(0)).z), _5.Load4(16).w > uint(int4(_5.Load4(0)).w));
bool4 _38 = bool4(_5.Load4(16).x >= uint(int4(_5.Load4(0)).x), _5.Load4(16).y >= uint(int4(_5.Load4(0)).y), _5.Load4(16).z >= uint(int4(_5.Load4(0)).z), _5.Load4(16).w >= uint(int4(_5.Load4(0)).w));
_6.Store4(0, uint4(_31.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _31.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _31.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _31.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
_6.Store4(0, uint4(_32.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _32.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _32.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _32.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
_6.Store4(0, uint4(_33.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _33.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _33.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _33.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
_6.Store4(0, uint4(_34.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _34.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _34.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _34.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
_6.Store4(0, uint4(_35.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _35.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _35.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _35.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
_6.Store4(0, uint4(_36.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _36.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _36.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _36.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
_6.Store4(0, uint4(_37.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _37.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _37.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _37.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
_6.Store4(0, uint4(_38.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _38.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _38.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _38.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
}
[numthreads(1, 1, 1)]
void main()
{
comp_main();
}

View File

@ -0,0 +1,29 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct _3
{
int4 _m0;
uint4 _m1;
};
struct _4
{
uint4 _m0;
int4 _m1;
};
kernel void main0(device _3& restrict _5 [[buffer(0)]], device _4& restrict _6 [[buffer(1)]])
{
_6._m0 = select(uint4(0u), uint4(1u), int4(_5._m1) < _5._m0);
_6._m0 = select(uint4(0u), uint4(1u), int4(_5._m1) <= _5._m0);
_6._m0 = select(uint4(0u), uint4(1u), _5._m1 < uint4(_5._m0));
_6._m0 = select(uint4(0u), uint4(1u), _5._m1 <= uint4(_5._m0));
_6._m0 = select(uint4(0u), uint4(1u), int4(_5._m1) > _5._m0);
_6._m0 = select(uint4(0u), uint4(1u), int4(_5._m1) >= _5._m0);
_6._m0 = select(uint4(0u), uint4(1u), _5._m1 > uint4(_5._m0));
_6._m0 = select(uint4(0u), uint4(1u), _5._m1 >= uint4(_5._m0));
}

View File

@ -0,0 +1,27 @@
#version 310 es
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(binding = 0, std430) restrict buffer _3_5
{
ivec4 _m0;
uvec4 _m1;
} _5;
layout(binding = 1, std430) restrict buffer _4_6
{
uvec4 _m0;
ivec4 _m1;
} _6;
void main()
{
_6._m0 = mix(uvec4(0u), uvec4(1u), lessThan(ivec4(_5._m1), _5._m0));
_6._m0 = mix(uvec4(0u), uvec4(1u), lessThanEqual(ivec4(_5._m1), _5._m0));
_6._m0 = mix(uvec4(0u), uvec4(1u), lessThan(_5._m1, uvec4(_5._m0)));
_6._m0 = mix(uvec4(0u), uvec4(1u), lessThanEqual(_5._m1, uvec4(_5._m0)));
_6._m0 = mix(uvec4(0u), uvec4(1u), greaterThan(ivec4(_5._m1), _5._m0));
_6._m0 = mix(uvec4(0u), uvec4(1u), greaterThanEqual(ivec4(_5._m1), _5._m0));
_6._m0 = mix(uvec4(0u), uvec4(1u), greaterThan(_5._m1, uvec4(_5._m0)));
_6._m0 = mix(uvec4(0u), uvec4(1u), greaterThanEqual(_5._m1, uvec4(_5._m0)));
}

View File

@ -0,0 +1,101 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 1
; Bound: 30
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %func "main"
OpExecutionMode %func LocalSize 1 1 1
OpSource ESSL 310
OpSourceExtension "GL_GOOGLE_cpp_style_line_directive"
OpSourceExtension "GL_GOOGLE_include_directive"
OpMemberDecorate %input_struct 0 Offset 0
OpMemberDecorate %input_struct 1 Offset 16
OpMemberDecorate %output_struct 0 Offset 0
OpMemberDecorate %output_struct 1 Offset 16
OpDecorate %input_struct BufferBlock
OpDecorate %inputs DescriptorSet 0
OpDecorate %inputs Binding 0
OpDecorate %inputs Restrict
OpDecorate %output_struct BufferBlock
OpDecorate %outputs DescriptorSet 0
OpDecorate %outputs Binding 1
OpDecorate %outputs Restrict
%void = OpTypeVoid
%main_func = OpTypeFunction %void
%bool = OpTypeBool
%bvec4 = OpTypeVector %bool 4
%uint = OpTypeInt 32 0
%uvec4 = OpTypeVector %uint 4
%int = OpTypeInt 32 1
%ivec4 = OpTypeVector %int 4
%ivec4_ptr = OpTypePointer Uniform %ivec4
%uvec4_ptr = OpTypePointer Uniform %uvec4
%zero = OpConstant %int 0
%one = OpConstant %int 1
%uzero = OpConstant %uint 0
%uone = OpConstant %uint 1
%utrue = OpConstantComposite %uvec4 %uone %uone %uone %uone
%ufalse = OpConstantComposite %uvec4 %uzero %uzero %uzero %uzero
%input_struct = OpTypeStruct %ivec4 %uvec4
%input_struct_ptr = OpTypePointer Uniform %input_struct
%inputs = OpVariable %input_struct_ptr Uniform
%output_struct = OpTypeStruct %uvec4 %ivec4
%output_struct_ptr = OpTypePointer Uniform %output_struct
%outputs = OpVariable %output_struct_ptr Uniform
%func = OpFunction %void None %main_func
%block = OpLabel
%input1_ptr = OpAccessChain %ivec4_ptr %inputs %zero
%input0_ptr = OpAccessChain %uvec4_ptr %inputs %one
%input1 = OpLoad %ivec4 %input1_ptr
%input0 = OpLoad %uvec4 %input0_ptr
%output_ptr_uvec4 = OpAccessChain %uvec4_ptr %outputs %zero
%result_slt = OpSLessThan %bvec4 %input0 %input1
%result_sle = OpSLessThanEqual %bvec4 %input0 %input1
%result_ult = OpULessThan %bvec4 %input0 %input1
%result_ule = OpULessThanEqual %bvec4 %input0 %input1
%result_sgt = OpSGreaterThan %bvec4 %input0 %input1
%result_sge = OpSGreaterThanEqual %bvec4 %input0 %input1
%result_ugt = OpUGreaterThan %bvec4 %input0 %input1
%result_uge = OpUGreaterThanEqual %bvec4 %input0 %input1
%int_slt = OpSelect %uvec4 %result_slt %utrue %ufalse
OpStore %output_ptr_uvec4 %int_slt
%int_sle = OpSelect %uvec4 %result_sle %utrue %ufalse
OpStore %output_ptr_uvec4 %int_sle
%int_ult = OpSelect %uvec4 %result_ult %utrue %ufalse
OpStore %output_ptr_uvec4 %int_ult
%int_ule = OpSelect %uvec4 %result_ule %utrue %ufalse
OpStore %output_ptr_uvec4 %int_ule
%int_sgt = OpSelect %uvec4 %result_sgt %utrue %ufalse
OpStore %output_ptr_uvec4 %int_sgt
%int_sge = OpSelect %uvec4 %result_sge %utrue %ufalse
OpStore %output_ptr_uvec4 %int_sge
%int_ugt = OpSelect %uvec4 %result_ugt %utrue %ufalse
OpStore %output_ptr_uvec4 %int_ugt
%int_uge = OpSelect %uvec4 %result_uge %utrue %ufalse
OpStore %output_ptr_uvec4 %int_uge
OpReturn
OpFunctionEnd

View File

@ -0,0 +1,101 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 1
; Bound: 30
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %func "main"
OpExecutionMode %func LocalSize 1 1 1
OpSource ESSL 310
OpSourceExtension "GL_GOOGLE_cpp_style_line_directive"
OpSourceExtension "GL_GOOGLE_include_directive"
OpMemberDecorate %input_struct 0 Offset 0
OpMemberDecorate %input_struct 1 Offset 16
OpMemberDecorate %output_struct 0 Offset 0
OpMemberDecorate %output_struct 1 Offset 16
OpDecorate %input_struct BufferBlock
OpDecorate %inputs DescriptorSet 0
OpDecorate %inputs Binding 0
OpDecorate %inputs Restrict
OpDecorate %output_struct BufferBlock
OpDecorate %outputs DescriptorSet 0
OpDecorate %outputs Binding 1
OpDecorate %outputs Restrict
%void = OpTypeVoid
%main_func = OpTypeFunction %void
%bool = OpTypeBool
%bvec4 = OpTypeVector %bool 4
%uint = OpTypeInt 32 0
%uvec4 = OpTypeVector %uint 4
%int = OpTypeInt 32 1
%ivec4 = OpTypeVector %int 4
%ivec4_ptr = OpTypePointer Uniform %ivec4
%uvec4_ptr = OpTypePointer Uniform %uvec4
%zero = OpConstant %int 0
%one = OpConstant %int 1
%uzero = OpConstant %uint 0
%uone = OpConstant %uint 1
%utrue = OpConstantComposite %uvec4 %uone %uone %uone %uone
%ufalse = OpConstantComposite %uvec4 %uzero %uzero %uzero %uzero
%input_struct = OpTypeStruct %ivec4 %uvec4
%input_struct_ptr = OpTypePointer Uniform %input_struct
%inputs = OpVariable %input_struct_ptr Uniform
%output_struct = OpTypeStruct %uvec4 %ivec4
%output_struct_ptr = OpTypePointer Uniform %output_struct
%outputs = OpVariable %output_struct_ptr Uniform
%func = OpFunction %void None %main_func
%block = OpLabel
%input1_ptr = OpAccessChain %ivec4_ptr %inputs %zero
%input0_ptr = OpAccessChain %uvec4_ptr %inputs %one
%input1 = OpLoad %ivec4 %input1_ptr
%input0 = OpLoad %uvec4 %input0_ptr
%output_ptr_uvec4 = OpAccessChain %uvec4_ptr %outputs %zero
%result_slt = OpSLessThan %bvec4 %input0 %input1
%result_sle = OpSLessThanEqual %bvec4 %input0 %input1
%result_ult = OpULessThan %bvec4 %input0 %input1
%result_ule = OpULessThanEqual %bvec4 %input0 %input1
%result_sgt = OpSGreaterThan %bvec4 %input0 %input1
%result_sge = OpSGreaterThanEqual %bvec4 %input0 %input1
%result_ugt = OpUGreaterThan %bvec4 %input0 %input1
%result_uge = OpUGreaterThanEqual %bvec4 %input0 %input1
%int_slt = OpSelect %uvec4 %result_slt %utrue %ufalse
OpStore %output_ptr_uvec4 %int_slt
%int_sle = OpSelect %uvec4 %result_sle %utrue %ufalse
OpStore %output_ptr_uvec4 %int_sle
%int_ult = OpSelect %uvec4 %result_ult %utrue %ufalse
OpStore %output_ptr_uvec4 %int_ult
%int_ule = OpSelect %uvec4 %result_ule %utrue %ufalse
OpStore %output_ptr_uvec4 %int_ule
%int_sgt = OpSelect %uvec4 %result_sgt %utrue %ufalse
OpStore %output_ptr_uvec4 %int_sgt
%int_sge = OpSelect %uvec4 %result_sge %utrue %ufalse
OpStore %output_ptr_uvec4 %int_sge
%int_ugt = OpSelect %uvec4 %result_ugt %utrue %ufalse
OpStore %output_ptr_uvec4 %int_ugt
%int_uge = OpSelect %uvec4 %result_uge %utrue %ufalse
OpStore %output_ptr_uvec4 %int_uge
OpReturn
OpFunctionEnd

View File

@ -0,0 +1,101 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 1
; Bound: 30
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %func "main"
OpExecutionMode %func LocalSize 1 1 1
OpSource ESSL 310
OpSourceExtension "GL_GOOGLE_cpp_style_line_directive"
OpSourceExtension "GL_GOOGLE_include_directive"
OpMemberDecorate %input_struct 0 Offset 0
OpMemberDecorate %input_struct 1 Offset 16
OpMemberDecorate %output_struct 0 Offset 0
OpMemberDecorate %output_struct 1 Offset 16
OpDecorate %input_struct BufferBlock
OpDecorate %inputs DescriptorSet 0
OpDecorate %inputs Binding 0
OpDecorate %inputs Restrict
OpDecorate %output_struct BufferBlock
OpDecorate %outputs DescriptorSet 0
OpDecorate %outputs Binding 1
OpDecorate %outputs Restrict
%void = OpTypeVoid
%main_func = OpTypeFunction %void
%bool = OpTypeBool
%bvec4 = OpTypeVector %bool 4
%uint = OpTypeInt 32 0
%uvec4 = OpTypeVector %uint 4
%int = OpTypeInt 32 1
%ivec4 = OpTypeVector %int 4
%ivec4_ptr = OpTypePointer Uniform %ivec4
%uvec4_ptr = OpTypePointer Uniform %uvec4
%zero = OpConstant %int 0
%one = OpConstant %int 1
%uzero = OpConstant %uint 0
%uone = OpConstant %uint 1
%utrue = OpConstantComposite %uvec4 %uone %uone %uone %uone
%ufalse = OpConstantComposite %uvec4 %uzero %uzero %uzero %uzero
%input_struct = OpTypeStruct %ivec4 %uvec4
%input_struct_ptr = OpTypePointer Uniform %input_struct
%inputs = OpVariable %input_struct_ptr Uniform
%output_struct = OpTypeStruct %uvec4 %ivec4
%output_struct_ptr = OpTypePointer Uniform %output_struct
%outputs = OpVariable %output_struct_ptr Uniform
%func = OpFunction %void None %main_func
%block = OpLabel
%input1_ptr = OpAccessChain %ivec4_ptr %inputs %zero
%input0_ptr = OpAccessChain %uvec4_ptr %inputs %one
%input1 = OpLoad %ivec4 %input1_ptr
%input0 = OpLoad %uvec4 %input0_ptr
%output_ptr_uvec4 = OpAccessChain %uvec4_ptr %outputs %zero
%result_slt = OpSLessThan %bvec4 %input0 %input1
%result_sle = OpSLessThanEqual %bvec4 %input0 %input1
%result_ult = OpULessThan %bvec4 %input0 %input1
%result_ule = OpULessThanEqual %bvec4 %input0 %input1
%result_sgt = OpSGreaterThan %bvec4 %input0 %input1
%result_sge = OpSGreaterThanEqual %bvec4 %input0 %input1
%result_ugt = OpUGreaterThan %bvec4 %input0 %input1
%result_uge = OpUGreaterThanEqual %bvec4 %input0 %input1
%int_slt = OpSelect %uvec4 %result_slt %utrue %ufalse
OpStore %output_ptr_uvec4 %int_slt
%int_sle = OpSelect %uvec4 %result_sle %utrue %ufalse
OpStore %output_ptr_uvec4 %int_sle
%int_ult = OpSelect %uvec4 %result_ult %utrue %ufalse
OpStore %output_ptr_uvec4 %int_ult
%int_ule = OpSelect %uvec4 %result_ule %utrue %ufalse
OpStore %output_ptr_uvec4 %int_ule
%int_sgt = OpSelect %uvec4 %result_sgt %utrue %ufalse
OpStore %output_ptr_uvec4 %int_sgt
%int_sge = OpSelect %uvec4 %result_sge %utrue %ufalse
OpStore %output_ptr_uvec4 %int_sge
%int_ugt = OpSelect %uvec4 %result_ugt %utrue %ufalse
OpStore %output_ptr_uvec4 %int_ugt
%int_uge = OpSelect %uvec4 %result_uge %utrue %ufalse
OpStore %output_ptr_uvec4 %int_uge
OpReturn
OpFunctionEnd

View File

@ -4159,8 +4159,18 @@ void CompilerGLSL::emit_unrolled_unary_op(uint32_t result_type, uint32_t result_
}
void CompilerGLSL::emit_unrolled_binary_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1,
const char *op, bool negate)
const char *op, bool negate, SPIRType::BaseType expected_type)
{
auto &type0 = expression_type(op0);
auto &type1 = expression_type(op1);
SPIRType target_type0 = type0;
SPIRType target_type1 = type1;
target_type0.basetype = expected_type;
target_type1.basetype = expected_type;
target_type0.vecsize = 1;
target_type1.vecsize = 1;
auto &type = get<SPIRType>(result_type);
auto expr = type_to_glsl_constructor(type);
expr += '(';
@ -4170,11 +4180,21 @@ void CompilerGLSL::emit_unrolled_binary_op(uint32_t result_type, uint32_t result
// that these expressions are properly flushed to temporaries if needed.
if (negate)
expr += "!(";
expr += to_extract_component_expression(op0, i);
if (expected_type != SPIRType::Unknown && type0.basetype != expected_type)
expr += bitcast_expression(target_type0, type0.basetype, to_extract_component_expression(op0, i));
else
expr += to_extract_component_expression(op0, i);
expr += ' ';
expr += op;
expr += ' ';
expr += to_extract_component_expression(op1, i);
if (expected_type != SPIRType::Unknown && type1.basetype != expected_type)
expr += bitcast_expression(target_type1, type1.basetype, to_extract_component_expression(op1, i));
else
expr += to_extract_component_expression(op1, i);
if (negate)
expr += ")";
@ -7820,6 +7840,10 @@ uint32_t CompilerGLSL::get_integer_width_for_instruction(const Instruction &inst
case OpSLessThanEqual:
case OpSGreaterThan:
case OpSGreaterThanEqual:
case OpULessThan:
case OpULessThanEqual:
case OpUGreaterThan:
case OpUGreaterThanEqual:
return expression_type(ops[2]).width;
default:
@ -8842,7 +8866,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
auto &type = get<SPIRType>(result_type);
if (type.vecsize > 1)
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "||", false);
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "||", false, SPIRType::Unknown);
else
GLSL_BOP(||);
break;
@ -8856,7 +8880,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
auto &type = get<SPIRType>(result_type);
if (type.vecsize > 1)
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "&&", false);
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "&&", false, SPIRType::Unknown);
else
GLSL_BOP(&&);
break;

View File

@ -471,7 +471,8 @@ protected:
void emit_unary_func_op(uint32_t result_type, uint32_t result_id, uint32_t op0, const char *op);
void emit_unrolled_unary_op(uint32_t result_type, uint32_t result_id, uint32_t operand, const char *op);
void emit_binary_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op);
void emit_unrolled_binary_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op, bool negate);
void emit_unrolled_binary_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op,
bool negate, SPIRType::BaseType expected_type);
void emit_binary_op_cast(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op,
SPIRType::BaseType input_type, bool skip_cast_if_equal_type);

View File

@ -4032,6 +4032,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
// If we need to do implicit bitcasts, make sure we do it with the correct type.
uint32_t integer_width = get_integer_width_for_instruction(instruction);
auto int_type = to_signed_basetype(integer_width);
auto uint_type = to_unsigned_basetype(integer_width);
switch (opcode)
{
@ -4194,7 +4195,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
auto id = ops[1];
if (expression_type(ops[2]).vecsize > 1)
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "==", false);
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "==", false, SPIRType::Unknown);
else
HLSL_BOP_CAST(==, int_type);
break;
@ -4214,7 +4215,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
auto id = ops[1];
if (expression_type(ops[2]).vecsize > 1)
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "==", false);
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "==", false, SPIRType::Unknown);
else
HLSL_BOP(==);
break;
@ -4226,7 +4227,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
auto id = ops[1];
if (expression_type(ops[2]).vecsize > 1)
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "!=", false);
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "!=", false, SPIRType::Unknown);
else
HLSL_BOP_CAST(!=, int_type);
break;
@ -4250,7 +4251,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
auto id = ops[1];
if (expression_type(ops[2]).vecsize > 1)
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "!=", false);
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "!=", false, SPIRType::Unknown);
else
HLSL_BOP(!=);
break;
@ -4261,10 +4262,10 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
{
auto result_type = ops[0];
auto id = ops[1];
auto type = opcode == OpUGreaterThan ? SPIRType::UInt : SPIRType::Int;
auto type = opcode == OpUGreaterThan ? uint_type : int_type;
if (expression_type(ops[2]).vecsize > 1)
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">", false);
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">", false, type);
else
HLSL_BOP_CAST(>, type);
break;
@ -4276,7 +4277,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
auto id = ops[1];
if (expression_type(ops[2]).vecsize > 1)
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">", false);
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">", false, SPIRType::Unknown);
else
HLSL_BOP(>);
break;
@ -4288,7 +4289,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
auto id = ops[1];
if (expression_type(ops[2]).vecsize > 1)
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<=", true);
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<=", true, SPIRType::Unknown);
else
CompilerGLSL::emit_instruction(instruction);
break;
@ -4300,9 +4301,9 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
auto result_type = ops[0];
auto id = ops[1];
auto type = opcode == OpUGreaterThanEqual ? SPIRType::UInt : SPIRType::Int;
auto type = opcode == OpUGreaterThanEqual ? uint_type : int_type;
if (expression_type(ops[2]).vecsize > 1)
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">=", false);
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">=", false, type);
else
HLSL_BOP_CAST(>=, type);
break;
@ -4314,7 +4315,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
auto id = ops[1];
if (expression_type(ops[2]).vecsize > 1)
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">=", false);
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">=", false, SPIRType::Unknown);
else
HLSL_BOP(>=);
break;
@ -4326,7 +4327,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
auto id = ops[1];
if (expression_type(ops[2]).vecsize > 1)
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<", true);
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<", true, SPIRType::Unknown);
else
CompilerGLSL::emit_instruction(instruction);
break;
@ -4338,9 +4339,9 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
auto result_type = ops[0];
auto id = ops[1];
auto type = opcode == OpULessThan ? SPIRType::UInt : SPIRType::Int;
auto type = opcode == OpULessThan ? uint_type : int_type;
if (expression_type(ops[2]).vecsize > 1)
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<", false);
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<", false, type);
else
HLSL_BOP_CAST(<, type);
break;
@ -4352,7 +4353,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
auto id = ops[1];
if (expression_type(ops[2]).vecsize > 1)
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<", false);
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<", false, SPIRType::Unknown);
else
HLSL_BOP(<);
break;
@ -4364,7 +4365,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
auto id = ops[1];
if (expression_type(ops[2]).vecsize > 1)
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">=", true);
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">=", true, SPIRType::Unknown);
else
CompilerGLSL::emit_instruction(instruction);
break;
@ -4376,9 +4377,9 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
auto result_type = ops[0];
auto id = ops[1];
auto type = opcode == OpULessThanEqual ? SPIRType::UInt : SPIRType::Int;
auto type = opcode == OpULessThanEqual ? uint_type : int_type;
if (expression_type(ops[2]).vecsize > 1)
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<=", false);
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<=", false, type);
else
HLSL_BOP_CAST(<=, type);
break;
@ -4390,7 +4391,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
auto id = ops[1];
if (expression_type(ops[2]).vecsize > 1)
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<=", false);
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<=", false, SPIRType::Unknown);
else
HLSL_BOP(<=);
break;
@ -4402,7 +4403,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
auto id = ops[1];
if (expression_type(ops[2]).vecsize > 1)
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">", true);
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">", true, SPIRType::Unknown);
else
CompilerGLSL::emit_instruction(instruction);
break;