diff options
| author | Andrzej Janik <[email protected]> | 2020-11-01 14:58:44 +0100 |
|---|---|---|
| committer | Andrzej Janik <[email protected]> | 2020-11-01 14:58:44 +0100 |
| commit | e5a53ed5d30fad3d8ebae6d72ead1564d2b97275 (patch) | |
| tree | 074d797ffdbafb530d11b592efefe84c277a5bea /ptx/src/test | |
| parent | b7d61baf37be52c7b2e5ea26be045470642f9a61 (diff) | |
| download | ZLUDA-e5a53ed5d30fad3d8ebae6d72ead1564d2b97275.tar.gz ZLUDA-e5a53ed5d30fad3d8ebae6d72ead1564d2b97275.zip | |
Implement neg instruction
Diffstat (limited to 'ptx/src/test')
| -rw-r--r-- | ptx/src/test/spirv_run/mod.rs | 1 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/neg.ptx | 21 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/neg.spvtxt | 47 |
3 files changed, 69 insertions, 0 deletions
diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index 4e9d39f..7ba3c4d 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/ptx/src/test/spirv_run/mod.rs @@ -104,6 +104,7 @@ test_ptx!(atom_add, [2u32, 4u32], [2u32, 6u32]); test_ptx!(div_approx, [1f32, 2f32], [0.5f32]);
test_ptx!(sqrt, [0.25f32], [0.5f32]);
test_ptx!(rsqrt, [0.25f64], [2f64]);
+test_ptx!(neg, [181i32], [-181i32]);
struct DisplayError<T: Debug> {
err: T,
diff --git a/ptx/src/test/spirv_run/neg.ptx b/ptx/src/test/spirv_run/neg.ptx new file mode 100644 index 0000000..60fe162 --- /dev/null +++ b/ptx/src/test/spirv_run/neg.ptx @@ -0,0 +1,21 @@ +.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry neg(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .reg .u64 in_addr;
+ .reg .u64 out_addr;
+ .reg .s32 temp1;
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+
+ ld.s32 temp1, [in_addr];
+ neg.s32 temp1, temp1;
+ st.s32 [out_addr], temp1;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/neg.spvtxt b/ptx/src/test/spirv_run/neg.spvtxt new file mode 100644 index 0000000..b358858 --- /dev/null +++ b/ptx/src/test/spirv_run/neg.spvtxt @@ -0,0 +1,47 @@ + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int64 + OpCapability Int8 + %26 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "not" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %29 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong +%_ptr_Generic_ulong = OpTypePointer Generic %ulong + %1 = OpFunction %void None %29 + %8 = OpFunctionParameter %ulong + %9 = OpFunctionParameter %ulong + %24 = 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_ulong Function + %7 = OpVariable %_ptr_Function_ulong Function + OpStore %2 %8 + OpStore %3 %9 + %11 = OpLoad %ulong %2 + %10 = OpCopyObject %ulong %11 + OpStore %4 %10 + %13 = OpLoad %ulong %3 + %12 = OpCopyObject %ulong %13 + OpStore %5 %12 + %15 = OpLoad %ulong %4 + %20 = OpConvertUToPtr %_ptr_Generic_ulong %15 + %14 = OpLoad %ulong %20 + OpStore %6 %14 + %17 = OpLoad %ulong %6 + %22 = OpCopyObject %ulong %17 + %21 = OpNot %ulong %22 + %16 = OpCopyObject %ulong %21 + OpStore %7 %16 + %18 = OpLoad %ulong %5 + %19 = OpLoad %ulong %7 + %23 = OpConvertUToPtr %_ptr_Generic_ulong %18 + OpStore %23 %19 + OpReturn + OpFunctionEnd |