diff options
Diffstat (limited to 'ptx/src/test')
| -rw-r--r-- | ptx/src/test/spirv_run/and.ptx | 23 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/and.spvtxt | 66 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/mod.rs | 1 |
3 files changed, 90 insertions, 0 deletions
diff --git a/ptx/src/test/spirv_run/and.ptx b/ptx/src/test/spirv_run/and.ptx new file mode 100644 index 0000000..88292a7 --- /dev/null +++ b/ptx/src/test/spirv_run/and.ptx @@ -0,0 +1,23 @@ +.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry and(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .reg .u64 in_addr;
+ .reg .u64 out_addr;
+ .reg .u32 temp1;
+ .reg .u32 temp2;
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+
+ ld.u32 temp1, [in_addr];
+ ld.u32 temp2, [in_addr+4];
+ and.b32 temp1, temp1, temp2;
+ st.u32 [out_addr], temp1;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/and.spvtxt b/ptx/src/test/spirv_run/and.spvtxt new file mode 100644 index 0000000..9b72477 --- /dev/null +++ b/ptx/src/test/spirv_run/and.spvtxt @@ -0,0 +1,66 @@ +; SPIR-V +; Version: 1.3 +; Generator: rspirv +; Bound: 41 +OpCapability GenericPointer +OpCapability Linkage +OpCapability Addresses +OpCapability Kernel +OpCapability Int8 +OpCapability Int16 +OpCapability Int64 +OpCapability Float16 +OpCapability Float64 +OpCapability FunctionFloatControlINTEL +OpExtension "SPV_INTEL_float_controls2" +%33 = OpExtInstImport "OpenCL.std" +OpMemoryModel Physical64 OpenCL +OpEntryPoint Kernel %1 "and" +%34 = OpTypeVoid +%35 = OpTypeInt 64 0 +%36 = OpTypeFunction %34 %35 %35 +%37 = OpTypePointer Function %35 +%38 = OpTypeInt 32 0 +%39 = OpTypePointer Function %38 +%40 = OpTypePointer Generic %38 +%23 = OpConstant %35 4 +%1 = OpFunction %34 None %36 +%8 = OpFunctionParameter %35 +%9 = OpFunctionParameter %35 +%31 = OpLabel +%2 = OpVariable %37 Function +%3 = OpVariable %37 Function +%4 = OpVariable %37 Function +%5 = OpVariable %37 Function +%6 = OpVariable %39 Function +%7 = OpVariable %39 Function +OpStore %2 %8 +OpStore %3 %9 +%11 = OpLoad %35 %2 +%10 = OpCopyObject %35 %11 +OpStore %4 %10 +%13 = OpLoad %35 %3 +%12 = OpCopyObject %35 %13 +OpStore %5 %12 +%15 = OpLoad %35 %4 +%25 = OpConvertUToPtr %40 %15 +%14 = OpLoad %38 %25 +OpStore %6 %14 +%17 = OpLoad %35 %4 +%24 = OpIAdd %35 %17 %23 +%26 = OpConvertUToPtr %40 %24 +%16 = OpLoad %38 %26 +OpStore %7 %16 +%19 = OpLoad %38 %6 +%20 = OpLoad %38 %7 +%28 = OpCopyObject %38 %19 +%29 = OpCopyObject %38 %20 +%27 = OpBitwiseAnd %38 %28 %29 +%18 = OpCopyObject %38 %27 +OpStore %6 %18 +%21 = OpLoad %35 %5 +%22 = OpLoad %38 %6 +%30 = OpConvertUToPtr %40 %21 +OpStore %30 %22 +OpReturn +OpFunctionEnd
\ No newline at end of file diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index 40acd46..dfdec72 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/ptx/src/test/spirv_run/mod.rs @@ -89,6 +89,7 @@ test_ptx!(rcp, [2f32], [0.5f32]); test_ptx!(mul_non_ftz, [0b1_00000000_10000000000000000000000u32, 0x3f000000u32], [0b1_00000000_01000000000000000000000u32]);
test_ptx!(constant_f32, [10f32], [5f32]);
test_ptx!(constant_negative, [-101i32], [101i32]);
+test_ptx!(and, [6u32, 3u32], [2u32]);
struct DisplayError<T: Debug> {
err: T,
|