From a6e6454d8bc0a4b75be7bfb738e281a1c9e4472c Mon Sep 17 00:00:00 2001 From: Andrzej Janik Date: Tue, 28 Jan 2025 01:46:34 +0000 Subject: [PATCH] Add failing test for ftz --- ptx/src/test/spirv_run/add_ftz.ptx | 24 ++++++++++++++++++++++++ ptx/src/test/spirv_run/mod.rs | 5 +++++ 2 files changed, 29 insertions(+) create mode 100644 ptx/src/test/spirv_run/add_ftz.ptx diff --git a/ptx/src/test/spirv_run/add_ftz.ptx b/ptx/src/test/spirv_run/add_ftz.ptx new file mode 100644 index 0000000..6909e96 --- /dev/null +++ b/ptx/src/test/spirv_run/add_ftz.ptx @@ -0,0 +1,24 @@ +.version 6.5 +.target sm_30 +.address_size 64 + +.visible .entry add_ftz( + .param .u64 input, + .param .u64 output +) +{ + .reg .u64 in_addr; + .reg .u64 out_addr; + .reg .f32 temp<4>; + + ld.param.u64 in_addr, [input]; + ld.param.u64 out_addr, [output]; + + ld.f32 temp0, [in_addr]; + ld.f32 temp1, [in_addr+4]; + add.ftz.f32 temp2, temp0, temp1; + add.f32 temp3, temp0, temp1; + st.f32 [out_addr], temp2; + st.f32 [out_addr+4], temp3; + ret; +} diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index e4171cd..99573f6 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/ptx/src/test/spirv_run/mod.rs @@ -180,6 +180,11 @@ test_ptx!(activemask, [0u32], [1u32]); test_ptx!(membar, [152731u32], [152731u32]); test_ptx!(shared_unify_extern, [7681u64, 7682u64], [15363u64]); test_ptx!(shared_unify_local, [16752u64, 714u64], [17466u64]); +test_ptx!( + add_ftz, + [f32::from_bits(0x800000), f32::from_bits(0x007FFFFF)], + [0x800000u32, 0xFFFFFF] +); test_ptx!(assertfail); test_ptx!(func_ptr);