Implement setp.nan and setp.num

This commit is contained in:
Andrzej Janik 2021-03-03 23:35:18 +01:00
parent 17291019e3
commit 15f465041d
8 changed files with 587 additions and 18 deletions

View file

@ -870,7 +870,7 @@ pub enum SetpCompareOp {
NanGreater,
NanGreaterOrEq,
IsNotNan,
IsNan,
IsAnyNan,
}
pub enum SetpBoolPostOp {

View file

@ -996,7 +996,7 @@ SetpCompareOp: ast::SetpCompareOp = {
".gtu" => ast::SetpCompareOp::NanGreater,
".geu" => ast::SetpCompareOp::NanGreaterOrEq,
".num" => ast::SetpCompareOp::IsNotNan,
".nan" => ast::SetpCompareOp::IsNan,
".nan" => ast::SetpCompareOp::IsAnyNan,
};
SetpBoolPostOp: ast::SetpBoolPostOp = {

View file

@ -139,11 +139,7 @@ test_ptx!(
[0b11111000_11000001_00100010_10100000u32, 16u32, 8u32],
[0b11000001u32]
);
test_ptx!(
bfi,
[0b10u32, 0b101u32, 0u32, 2u32],
[0b110u32]
);
test_ptx!(bfi, [0b10u32, 0b101u32, 0u32, 2u32], [0b110u32]);
test_ptx!(stateful_ld_st_simple, [121u64], [121u64]);
test_ptx!(stateful_ld_st_ntid, [123u64], [123u64]);
test_ptx!(stateful_ld_st_ntid_chain, [12651u64], [12651u64]);
@ -156,6 +152,34 @@ test_ptx!(add_tuning, [2u64], [3u64]);
test_ptx!(add_non_coherent, [3u64], [4u64]);
test_ptx!(sign_extend, [-1i16], [-1i32]);
test_ptx!(atom_add_float, [1.25f32, 0.5f32], [1.25f32, 1.75f32]);
test_ptx!(
setp_nan,
[
0.5f32,
f32::NAN,
f32::NAN,
0.5f32,
f32::NAN,
f32::NAN,
0.5f32,
0.5f32
],
[1u32, 1u32, 1u32, 0u32]
);
test_ptx!(
setp_num,
[
0.5f32,
f32::NAN,
f32::NAN,
0.5f32,
f32::NAN,
f32::NAN,
0.5f32,
0.5f32
],
[0u32, 0u32, 0u32, 2u32]
);
struct DisplayError<T: Debug> {
err: T,

View file

@ -0,0 +1,51 @@
.version 6.5
.target sm_30
.address_size 64
.visible .entry setp_nan(
.param .u64 input,
.param .u64 output
)
{
.reg .u64 in_addr;
.reg .u64 out_addr;
.reg .f32 pair1_1;
.reg .f32 pair1_2;
.reg .f32 pair2_1;
.reg .f32 pair2_2;
.reg .f32 pair3_1;
.reg .f32 pair3_2;
.reg .f32 pair4_1;
.reg .f32 pair4_2;
.reg .u32 temp;
.reg .pred pred;
ld.param.u64 in_addr, [input];
ld.param.u64 out_addr, [output];
ld.f32 pair1_1, [in_addr];
ld.f32 pair1_2, [in_addr + 4];
ld.f32 pair2_1, [in_addr + 8];
ld.f32 pair2_2, [in_addr + 12];
ld.f32 pair3_1, [in_addr + 16];
ld.f32 pair3_2, [in_addr + 20];
ld.f32 pair4_1, [in_addr + 24];
ld.f32 pair4_2, [in_addr + 28];
setp.nan.f32 pred, pair1_1, pair1_2;
@pred mov.u32 temp, 1;
@!pred mov.u32 temp, 0;
st.u32 [out_addr], temp;
setp.nan.f32 pred, pair2_1, pair2_2;
@pred mov.u32 temp, 1;
@!pred mov.u32 temp, 0;
st.u32 [out_addr + 4], temp;
setp.nan.f32 pred, pair3_1, pair3_2;
@pred mov.u32 temp, 1;
@!pred mov.u32 temp, 0;
st.u32 [out_addr + 8], temp;
setp.nan.f32 pred, pair4_1, pair4_2;
@pred mov.u32 temp, 1;
@!pred mov.u32 temp, 0;
st.u32 [out_addr + 12], temp;
ret;
}

View file

@ -0,0 +1,206 @@
OpCapability GenericPointer
OpCapability Linkage
OpCapability Addresses
OpCapability Kernel
OpCapability Int8
OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%130 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "setp_nan"
%void = OpTypeVoid
%ulong = OpTypeInt 64 0
%133 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong
%float = OpTypeFloat 32
%_ptr_Function_float = OpTypePointer Function %float
%uint = OpTypeInt 32 0
%_ptr_Function_uint = OpTypePointer Function %uint
%bool = OpTypeBool
%_ptr_Function_bool = OpTypePointer Function %bool
%_ptr_Generic_float = OpTypePointer Generic %float
%ulong_4 = OpConstant %ulong 4
%ulong_8 = OpConstant %ulong 8
%ulong_12 = OpConstant %ulong 12
%ulong_16 = OpConstant %ulong 16
%ulong_20 = OpConstant %ulong 20
%ulong_24 = OpConstant %ulong 24
%ulong_28 = OpConstant %ulong 28
%uint_1 = OpConstant %uint 1
%uint_0 = OpConstant %uint 0
%_ptr_Generic_uint = OpTypePointer Generic %uint
%uint_1_0 = OpConstant %uint 1
%uint_0_0 = OpConstant %uint 0
%ulong_4_0 = OpConstant %ulong 4
%uint_1_1 = OpConstant %uint 1
%uint_0_1 = OpConstant %uint 0
%ulong_8_0 = OpConstant %ulong 8
%uint_1_2 = OpConstant %uint 1
%uint_0_2 = OpConstant %uint 0
%ulong_12_0 = OpConstant %ulong 12
%1 = OpFunction %void None %133
%32 = OpFunctionParameter %ulong
%33 = OpFunctionParameter %ulong
%128 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function
%5 = OpVariable %_ptr_Function_ulong Function
%6 = OpVariable %_ptr_Function_float Function
%7 = OpVariable %_ptr_Function_float Function
%8 = OpVariable %_ptr_Function_float Function
%9 = OpVariable %_ptr_Function_float Function
%10 = OpVariable %_ptr_Function_float Function
%11 = OpVariable %_ptr_Function_float Function
%12 = OpVariable %_ptr_Function_float Function
%13 = OpVariable %_ptr_Function_float Function
%14 = OpVariable %_ptr_Function_uint Function
%15 = OpVariable %_ptr_Function_bool Function
OpStore %2 %32
OpStore %3 %33
%34 = OpLoad %ulong %2 Aligned 8
OpStore %4 %34
%35 = OpLoad %ulong %3 Aligned 8
OpStore %5 %35
%37 = OpLoad %ulong %4
%116 = OpConvertUToPtr %_ptr_Generic_float %37
%36 = OpLoad %float %116 Aligned 4
OpStore %6 %36
%39 = OpLoad %ulong %4
%89 = OpIAdd %ulong %39 %ulong_4
%117 = OpConvertUToPtr %_ptr_Generic_float %89
%38 = OpLoad %float %117 Aligned 4
OpStore %7 %38
%41 = OpLoad %ulong %4
%91 = OpIAdd %ulong %41 %ulong_8
%118 = OpConvertUToPtr %_ptr_Generic_float %91
%40 = OpLoad %float %118 Aligned 4
OpStore %8 %40
%43 = OpLoad %ulong %4
%93 = OpIAdd %ulong %43 %ulong_12
%119 = OpConvertUToPtr %_ptr_Generic_float %93
%42 = OpLoad %float %119 Aligned 4
OpStore %9 %42
%45 = OpLoad %ulong %4
%95 = OpIAdd %ulong %45 %ulong_16
%120 = OpConvertUToPtr %_ptr_Generic_float %95
%44 = OpLoad %float %120 Aligned 4
OpStore %10 %44
%47 = OpLoad %ulong %4
%97 = OpIAdd %ulong %47 %ulong_20
%121 = OpConvertUToPtr %_ptr_Generic_float %97
%46 = OpLoad %float %121 Aligned 4
OpStore %11 %46
%49 = OpLoad %ulong %4
%99 = OpIAdd %ulong %49 %ulong_24
%122 = OpConvertUToPtr %_ptr_Generic_float %99
%48 = OpLoad %float %122 Aligned 4
OpStore %12 %48
%51 = OpLoad %ulong %4
%101 = OpIAdd %ulong %51 %ulong_28
%123 = OpConvertUToPtr %_ptr_Generic_float %101
%50 = OpLoad %float %123 Aligned 4
OpStore %13 %50
%53 = OpLoad %float %6
%54 = OpLoad %float %7
%142 = OpIsNan %bool %53
%143 = OpIsNan %bool %54
%52 = OpLogicalOr %bool %142 %143
OpStore %15 %52
%55 = OpLoad %bool %15
OpBranchConditional %55 %16 %17
%16 = OpLabel
%56 = OpCopyObject %uint %uint_1
OpStore %14 %56
OpBranch %17
%17 = OpLabel
%57 = OpLoad %bool %15
OpBranchConditional %57 %19 %18
%18 = OpLabel
%58 = OpCopyObject %uint %uint_0
OpStore %14 %58
OpBranch %19
%19 = OpLabel
%59 = OpLoad %ulong %5
%60 = OpLoad %uint %14
%124 = OpConvertUToPtr %_ptr_Generic_uint %59
OpStore %124 %60 Aligned 4
%62 = OpLoad %float %8
%63 = OpLoad %float %9
%145 = OpIsNan %bool %62
%146 = OpIsNan %bool %63
%61 = OpLogicalOr %bool %145 %146
OpStore %15 %61
%64 = OpLoad %bool %15
OpBranchConditional %64 %20 %21
%20 = OpLabel
%65 = OpCopyObject %uint %uint_1_0
OpStore %14 %65
OpBranch %21
%21 = OpLabel
%66 = OpLoad %bool %15
OpBranchConditional %66 %23 %22
%22 = OpLabel
%67 = OpCopyObject %uint %uint_0_0
OpStore %14 %67
OpBranch %23
%23 = OpLabel
%68 = OpLoad %ulong %5
%69 = OpLoad %uint %14
%107 = OpIAdd %ulong %68 %ulong_4_0
%125 = OpConvertUToPtr %_ptr_Generic_uint %107
OpStore %125 %69 Aligned 4
%71 = OpLoad %float %10
%72 = OpLoad %float %11
%147 = OpIsNan %bool %71
%148 = OpIsNan %bool %72
%70 = OpLogicalOr %bool %147 %148
OpStore %15 %70
%73 = OpLoad %bool %15
OpBranchConditional %73 %24 %25
%24 = OpLabel
%74 = OpCopyObject %uint %uint_1_1
OpStore %14 %74
OpBranch %25
%25 = OpLabel
%75 = OpLoad %bool %15
OpBranchConditional %75 %27 %26
%26 = OpLabel
%76 = OpCopyObject %uint %uint_0_1
OpStore %14 %76
OpBranch %27
%27 = OpLabel
%77 = OpLoad %ulong %5
%78 = OpLoad %uint %14
%111 = OpIAdd %ulong %77 %ulong_8_0
%126 = OpConvertUToPtr %_ptr_Generic_uint %111
OpStore %126 %78 Aligned 4
%80 = OpLoad %float %12
%81 = OpLoad %float %13
%149 = OpIsNan %bool %80
%150 = OpIsNan %bool %81
%79 = OpLogicalOr %bool %149 %150
OpStore %15 %79
%82 = OpLoad %bool %15
OpBranchConditional %82 %28 %29
%28 = OpLabel
%83 = OpCopyObject %uint %uint_1_2
OpStore %14 %83
OpBranch %29
%29 = OpLabel
%84 = OpLoad %bool %15
OpBranchConditional %84 %31 %30
%30 = OpLabel
%85 = OpCopyObject %uint %uint_0_2
OpStore %14 %85
OpBranch %31
%31 = OpLabel
%86 = OpLoad %ulong %5
%87 = OpLoad %uint %14
%115 = OpIAdd %ulong %86 %ulong_12_0
%127 = OpConvertUToPtr %_ptr_Generic_uint %115
OpStore %127 %87 Aligned 4
OpReturn
OpFunctionEnd

View file

@ -0,0 +1,51 @@
.version 6.5
.target sm_30
.address_size 64
.visible .entry setp_num(
.param .u64 input,
.param .u64 output
)
{
.reg .u64 in_addr;
.reg .u64 out_addr;
.reg .f32 pair1_1;
.reg .f32 pair1_2;
.reg .f32 pair2_1;
.reg .f32 pair2_2;
.reg .f32 pair3_1;
.reg .f32 pair3_2;
.reg .f32 pair4_1;
.reg .f32 pair4_2;
.reg .u32 temp;
.reg .pred pred;
ld.param.u64 in_addr, [input];
ld.param.u64 out_addr, [output];
ld.f32 pair1_1, [in_addr];
ld.f32 pair1_2, [in_addr + 4];
ld.f32 pair2_1, [in_addr + 8];
ld.f32 pair2_2, [in_addr + 12];
ld.f32 pair3_1, [in_addr + 16];
ld.f32 pair3_2, [in_addr + 20];
ld.f32 pair4_1, [in_addr + 24];
ld.f32 pair4_2, [in_addr + 28];
setp.num.f32 pred, pair1_1, pair1_2;
@pred mov.u32 temp, 2;
@!pred mov.u32 temp, 0;
st.u32 [out_addr], temp;
setp.num.f32 pred, pair2_1, pair2_2;
@pred mov.u32 temp, 2;
@!pred mov.u32 temp, 0;
st.u32 [out_addr + 4], temp;
setp.num.f32 pred, pair3_1, pair3_2;
@pred mov.u32 temp, 2;
@!pred mov.u32 temp, 0;
st.u32 [out_addr + 8], temp;
setp.num.f32 pred, pair4_1, pair4_2;
@pred mov.u32 temp, 2;
@!pred mov.u32 temp, 0;
st.u32 [out_addr + 12], temp;
ret;
}

View file

@ -0,0 +1,218 @@
OpCapability GenericPointer
OpCapability Linkage
OpCapability Addresses
OpCapability Kernel
OpCapability Int8
OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%130 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "setp_num"
%void = OpTypeVoid
%ulong = OpTypeInt 64 0
%133 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong
%float = OpTypeFloat 32
%_ptr_Function_float = OpTypePointer Function %float
%uint = OpTypeInt 32 0
%_ptr_Function_uint = OpTypePointer Function %uint
%bool = OpTypeBool
%_ptr_Function_bool = OpTypePointer Function %bool
%_ptr_Generic_float = OpTypePointer Generic %float
%ulong_4 = OpConstant %ulong 4
%ulong_8 = OpConstant %ulong 8
%ulong_12 = OpConstant %ulong 12
%ulong_16 = OpConstant %ulong 16
%ulong_20 = OpConstant %ulong 20
%ulong_24 = OpConstant %ulong 24
%ulong_28 = OpConstant %ulong 28
%true = OpConstantTrue %bool
%false = OpConstantFalse %bool
%uint_2 = OpConstant %uint 2
%uint_0 = OpConstant %uint 0
%_ptr_Generic_uint = OpTypePointer Generic %uint
%true_0 = OpConstantTrue %bool
%false_0 = OpConstantFalse %bool
%uint_2_0 = OpConstant %uint 2
%uint_0_0 = OpConstant %uint 0
%ulong_4_0 = OpConstant %ulong 4
%true_1 = OpConstantTrue %bool
%false_1 = OpConstantFalse %bool
%uint_2_1 = OpConstant %uint 2
%uint_0_1 = OpConstant %uint 0
%ulong_8_0 = OpConstant %ulong 8
%true_2 = OpConstantTrue %bool
%false_2 = OpConstantFalse %bool
%uint_2_2 = OpConstant %uint 2
%uint_0_2 = OpConstant %uint 0
%ulong_12_0 = OpConstant %ulong 12
%1 = OpFunction %void None %133
%32 = OpFunctionParameter %ulong
%33 = OpFunctionParameter %ulong
%128 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function
%5 = OpVariable %_ptr_Function_ulong Function
%6 = OpVariable %_ptr_Function_float Function
%7 = OpVariable %_ptr_Function_float Function
%8 = OpVariable %_ptr_Function_float Function
%9 = OpVariable %_ptr_Function_float Function
%10 = OpVariable %_ptr_Function_float Function
%11 = OpVariable %_ptr_Function_float Function
%12 = OpVariable %_ptr_Function_float Function
%13 = OpVariable %_ptr_Function_float Function
%14 = OpVariable %_ptr_Function_uint Function
%15 = OpVariable %_ptr_Function_bool Function
OpStore %2 %32
OpStore %3 %33
%34 = OpLoad %ulong %2 Aligned 8
OpStore %4 %34
%35 = OpLoad %ulong %3 Aligned 8
OpStore %5 %35
%37 = OpLoad %ulong %4
%116 = OpConvertUToPtr %_ptr_Generic_float %37
%36 = OpLoad %float %116 Aligned 4
OpStore %6 %36
%39 = OpLoad %ulong %4
%89 = OpIAdd %ulong %39 %ulong_4
%117 = OpConvertUToPtr %_ptr_Generic_float %89
%38 = OpLoad %float %117 Aligned 4
OpStore %7 %38
%41 = OpLoad %ulong %4
%91 = OpIAdd %ulong %41 %ulong_8
%118 = OpConvertUToPtr %_ptr_Generic_float %91
%40 = OpLoad %float %118 Aligned 4
OpStore %8 %40
%43 = OpLoad %ulong %4
%93 = OpIAdd %ulong %43 %ulong_12
%119 = OpConvertUToPtr %_ptr_Generic_float %93
%42 = OpLoad %float %119 Aligned 4
OpStore %9 %42
%45 = OpLoad %ulong %4
%95 = OpIAdd %ulong %45 %ulong_16
%120 = OpConvertUToPtr %_ptr_Generic_float %95
%44 = OpLoad %float %120 Aligned 4
OpStore %10 %44
%47 = OpLoad %ulong %4
%97 = OpIAdd %ulong %47 %ulong_20
%121 = OpConvertUToPtr %_ptr_Generic_float %97
%46 = OpLoad %float %121 Aligned 4
OpStore %11 %46
%49 = OpLoad %ulong %4
%99 = OpIAdd %ulong %49 %ulong_24
%122 = OpConvertUToPtr %_ptr_Generic_float %99
%48 = OpLoad %float %122 Aligned 4
OpStore %12 %48
%51 = OpLoad %ulong %4
%101 = OpIAdd %ulong %51 %ulong_28
%123 = OpConvertUToPtr %_ptr_Generic_float %101
%50 = OpLoad %float %123 Aligned 4
OpStore %13 %50
%53 = OpLoad %float %6
%54 = OpLoad %float %7
%142 = OpIsNan %bool %53
%143 = OpIsNan %bool %54
%144 = OpLogicalOr %bool %142 %143
%52 = OpSelect %bool %144 %false %true
OpStore %15 %52
%55 = OpLoad %bool %15
OpBranchConditional %55 %16 %17
%16 = OpLabel
%56 = OpCopyObject %uint %uint_2
OpStore %14 %56
OpBranch %17
%17 = OpLabel
%57 = OpLoad %bool %15
OpBranchConditional %57 %19 %18
%18 = OpLabel
%58 = OpCopyObject %uint %uint_0
OpStore %14 %58
OpBranch %19
%19 = OpLabel
%59 = OpLoad %ulong %5
%60 = OpLoad %uint %14
%124 = OpConvertUToPtr %_ptr_Generic_uint %59
OpStore %124 %60 Aligned 4
%62 = OpLoad %float %8
%63 = OpLoad %float %9
%148 = OpIsNan %bool %62
%149 = OpIsNan %bool %63
%150 = OpLogicalOr %bool %148 %149
%61 = OpSelect %bool %150 %false_0 %true_0
OpStore %15 %61
%64 = OpLoad %bool %15
OpBranchConditional %64 %20 %21
%20 = OpLabel
%65 = OpCopyObject %uint %uint_2_0
OpStore %14 %65
OpBranch %21
%21 = OpLabel
%66 = OpLoad %bool %15
OpBranchConditional %66 %23 %22
%22 = OpLabel
%67 = OpCopyObject %uint %uint_0_0
OpStore %14 %67
OpBranch %23
%23 = OpLabel
%68 = OpLoad %ulong %5
%69 = OpLoad %uint %14
%107 = OpIAdd %ulong %68 %ulong_4_0
%125 = OpConvertUToPtr %_ptr_Generic_uint %107
OpStore %125 %69 Aligned 4
%71 = OpLoad %float %10
%72 = OpLoad %float %11
%153 = OpIsNan %bool %71
%154 = OpIsNan %bool %72
%155 = OpLogicalOr %bool %153 %154
%70 = OpSelect %bool %155 %false_1 %true_1
OpStore %15 %70
%73 = OpLoad %bool %15
OpBranchConditional %73 %24 %25
%24 = OpLabel
%74 = OpCopyObject %uint %uint_2_1
OpStore %14 %74
OpBranch %25
%25 = OpLabel
%75 = OpLoad %bool %15
OpBranchConditional %75 %27 %26
%26 = OpLabel
%76 = OpCopyObject %uint %uint_0_1
OpStore %14 %76
OpBranch %27
%27 = OpLabel
%77 = OpLoad %ulong %5
%78 = OpLoad %uint %14
%111 = OpIAdd %ulong %77 %ulong_8_0
%126 = OpConvertUToPtr %_ptr_Generic_uint %111
OpStore %126 %78 Aligned 4
%80 = OpLoad %float %12
%81 = OpLoad %float %13
%158 = OpIsNan %bool %80
%159 = OpIsNan %bool %81
%160 = OpLogicalOr %bool %158 %159
%79 = OpSelect %bool %160 %false_2 %true_2
OpStore %15 %79
%82 = OpLoad %bool %15
OpBranchConditional %82 %28 %29
%28 = OpLabel
%83 = OpCopyObject %uint %uint_2_2
OpStore %14 %83
OpBranch %29
%29 = OpLabel
%84 = OpLoad %bool %15
OpBranchConditional %84 %31 %30
%30 = OpLabel
%85 = OpCopyObject %uint %uint_0_2
OpStore %14 %85
OpBranch %31
%31 = OpLabel
%86 = OpLoad %ulong %5
%87 = OpLoad %uint %14
%115 = OpIAdd %ulong %86 %ulong_12_0
%127 = OpConvertUToPtr %_ptr_Generic_uint %115
OpStore %127 %87 Aligned 4
OpReturn
OpFunctionEnd

View file

@ -3031,16 +3031,7 @@ fn emit_function_body_ops(
let operand = a.src;
match t {
ast::BooleanType::Pred => {
// HACK ALERT
// Temporary workaround until IGC gets its shit together
// Currently IGC carries two copies of SPIRV-LLVM translator
// a new one in /llvm-spirv/ and old one in /IGC/AdaptorOCL/SPIRV/.
// Obviously, old and buggy one is used for compiling L0 SPIRV
// https://github.com/intel/intel-graphics-compiler/issues/148
let type_pred = map.get_or_add_scalar(builder, ast::ScalarType::Pred);
let const_true = builder.constant_true(type_pred, None);
let const_false = builder.constant_false(type_pred, None);
builder.select(result_type, result_id, operand, const_false, const_true)
logical_not(builder, result_type, result_id, operand)
}
_ => builder.not(result_type, result_id, operand),
}?;
@ -3426,7 +3417,7 @@ fn emit_logical_xor_spirv(
) -> Result<spirv::Word, dr::Error> {
let temp_or = builder.logical_or(result_type, None, op1, op2)?;
let temp_and = builder.logical_and(result_type, None, op1, op2)?;
let temp_neg = builder.logical_not(result_type, None, temp_and)?;
let temp_neg = logical_not(builder, result_type, None, temp_and)?;
builder.logical_and(result_type, result_id, temp_or, temp_neg)
}
@ -4072,11 +4063,39 @@ fn emit_setp(
(ast::SetpCompareOp::NanGreaterOrEq, _) => {
builder.f_unord_greater_than_equal(result_type, result_id, operand_1, operand_2)
}
(ast::SetpCompareOp::IsAnyNan, _) => {
let temp1 = builder.is_nan(result_type, None, operand_1)?;
let temp2 = builder.is_nan(result_type, None, operand_2)?;
builder.logical_or(result_type, result_id, temp1, temp2)
}
(ast::SetpCompareOp::IsNotNan, _) => {
let temp1 = builder.is_nan(result_type, None, operand_1)?;
let temp2 = builder.is_nan(result_type, None, operand_2)?;
let any_nan = builder.logical_or(result_type, None, temp1, temp2)?;
logical_not(builder, result_type, result_id, any_nan)
}
_ => todo!(),
}?;
Ok(())
}
// HACK ALERT
// Temporary workaround until IGC gets its shit together
// Currently IGC carries two copies of SPIRV-LLVM translator
// a new one in /llvm-spirv/ and old one in /IGC/AdaptorOCL/SPIRV/.
// Obviously, old and buggy one is used for compiling L0 SPIRV
// https://github.com/intel/intel-graphics-compiler/issues/148
fn logical_not(
builder: &mut dr::Builder,
result_type: spirv::Word,
result_id: Option<spirv::Word>,
operand: spirv::Word,
) -> Result<spirv::Word, dr::Error> {
let const_true = builder.constant_true(result_type, None);
let const_false = builder.constant_false(result_type, None);
builder.select(result_type, result_id, operand, const_false, const_true)
}
fn emit_mul_sint(
builder: &mut dr::Builder,
map: &mut TypeWordMap,