From 15f465041d2d19971455b70d778af51d50aca4fe Mon Sep 17 00:00:00 2001 From: Andrzej Janik Date: Wed, 3 Mar 2021 23:35:18 +0100 Subject: [PATCH] Implement setp.nan and setp.num --- ptx/src/ast.rs | 2 +- ptx/src/ptx.lalrpop | 2 +- ptx/src/test/spirv_run/mod.rs | 34 +++- ptx/src/test/spirv_run/setp_nan.ptx | 51 ++++++ ptx/src/test/spirv_run/setp_nan.spvtxt | 206 +++++++++++++++++++++++ ptx/src/test/spirv_run/setp_num.ptx | 51 ++++++ ptx/src/test/spirv_run/setp_num.spvtxt | 218 +++++++++++++++++++++++++ ptx/src/translate.rs | 41 +++-- 8 files changed, 587 insertions(+), 18 deletions(-) create mode 100644 ptx/src/test/spirv_run/setp_nan.ptx create mode 100644 ptx/src/test/spirv_run/setp_nan.spvtxt create mode 100644 ptx/src/test/spirv_run/setp_num.ptx create mode 100644 ptx/src/test/spirv_run/setp_num.spvtxt diff --git a/ptx/src/ast.rs b/ptx/src/ast.rs index dc9d4cc..1c6d2fb 100644 --- a/ptx/src/ast.rs +++ b/ptx/src/ast.rs @@ -870,7 +870,7 @@ pub enum SetpCompareOp { NanGreater, NanGreaterOrEq, IsNotNan, - IsNan, + IsAnyNan, } pub enum SetpBoolPostOp { diff --git a/ptx/src/ptx.lalrpop b/ptx/src/ptx.lalrpop index 6a0ca55..423fd57 100644 --- a/ptx/src/ptx.lalrpop +++ b/ptx/src/ptx.lalrpop @@ -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 = { diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index c802320..14d3284 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/ptx/src/test/spirv_run/mod.rs @@ -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 { err: T, diff --git a/ptx/src/test/spirv_run/setp_nan.ptx b/ptx/src/test/spirv_run/setp_nan.ptx new file mode 100644 index 0000000..6a9951e --- /dev/null +++ b/ptx/src/test/spirv_run/setp_nan.ptx @@ -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; +} diff --git a/ptx/src/test/spirv_run/setp_nan.spvtxt b/ptx/src/test/spirv_run/setp_nan.spvtxt new file mode 100644 index 0000000..4a9fe11 --- /dev/null +++ b/ptx/src/test/spirv_run/setp_nan.spvtxt @@ -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 diff --git a/ptx/src/test/spirv_run/setp_num.ptx b/ptx/src/test/spirv_run/setp_num.ptx new file mode 100644 index 0000000..d83ea4e --- /dev/null +++ b/ptx/src/test/spirv_run/setp_num.ptx @@ -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; +} diff --git a/ptx/src/test/spirv_run/setp_num.spvtxt b/ptx/src/test/spirv_run/setp_num.spvtxt new file mode 100644 index 0000000..3ac6eab --- /dev/null +++ b/ptx/src/test/spirv_run/setp_num.spvtxt @@ -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 diff --git a/ptx/src/translate.rs b/ptx/src/translate.rs index 3291ad5..e39280a 100644 --- a/ptx/src/translate.rs +++ b/ptx/src/translate.rs @@ -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 { 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, + operand: spirv::Word, +) -> Result { + 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,