diff options
| author | Andrzej Janik <[email protected]> | 2020-11-18 02:26:05 +0100 |
|---|---|---|
| committer | Andrzej Janik <[email protected]> | 2020-11-18 02:26:05 +0100 |
| commit | b0b0c21a9b5eb973d366dda6899e4a94442f923d (patch) | |
| tree | e1e8b0015c17c15d268a12329f7c3c6d1138c2af | |
| parent | ac45c2bde42cef74d68b572920b0749f8eea6837 (diff) | |
| download | ZLUDA-b0b0c21a9b5eb973d366dda6899e4a94442f923d.tar.gz ZLUDA-b0b0c21a9b5eb973d366dda6899e4a94442f923d.zip | |
Add more complicated test
| -rw-r--r-- | ptx/src/test/spirv_run/mod.rs | 2 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/stateful_ld_st_ntid.ptx | 31 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/stateful_ld_st_ntid.spvtxt | 57 |
3 files changed, 90 insertions, 0 deletions
diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index 0814dca..ab79a88 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/ptx/src/test/spirv_run/mod.rs @@ -134,6 +134,8 @@ test_ptx!( [0b11000001u32]
);
test_ptx!(stateful_ld_st_simple, [121u64], [121u64]);
+test_ptx!(stateful_ld_st_ntid, [123u64], [123u64]);
+
struct DisplayError<T: Debug> {
err: T,
diff --git a/ptx/src/test/spirv_run/stateful_ld_st_ntid.ptx b/ptx/src/test/spirv_run/stateful_ld_st_ntid.ptx new file mode 100644 index 0000000..1fc37d1 --- /dev/null +++ b/ptx/src/test/spirv_run/stateful_ld_st_ntid.ptx @@ -0,0 +1,31 @@ +.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry stateful_ld_st_ntid(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .reg .b64 in_addr;
+ .reg .b64 out_addr;
+ .reg .u32 tid_32;
+ .reg .u64 tid_64;
+ .reg .u64 temp;
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+
+ cvta.to.global.u64 in_addr, in_addr;
+ cvta.to.global.u64 out_addr, out_addr;
+
+ mov.u32 tid_32, %tid.x;
+ cvt.u64.u32 tid_64, tid_32;
+
+ add.u64 in_addr, in_addr, tid_64;
+ add.u64 out_addr, out_addr, tid_64;
+
+ ld.global.u64 temp, [in_addr];
+ st.global.u64 [out_addr], temp;
+ ret;
+}
\ No newline at end of file diff --git a/ptx/src/test/spirv_run/stateful_ld_st_ntid.spvtxt b/ptx/src/test/spirv_run/stateful_ld_st_ntid.spvtxt new file mode 100644 index 0000000..963d88a --- /dev/null +++ b/ptx/src/test/spirv_run/stateful_ld_st_ntid.spvtxt @@ -0,0 +1,57 @@ + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %30 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "ld_st_offset" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %33 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %uint = OpTypeInt 32 0 +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %ulong_4 = OpConstant %ulong 4 + %ulong_4_0 = OpConstant %ulong 4 + %1 = OpFunction %void None %33 + %8 = OpFunctionParameter %ulong + %9 = OpFunctionParameter %ulong + %28 = 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_uint Function + %7 = OpVariable %_ptr_Function_uint Function + OpStore %2 %8 + OpStore %3 %9 + %10 = OpLoad %ulong %2 + OpStore %4 %10 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %24 = OpConvertUToPtr %_ptr_Generic_uint %13 + %12 = OpLoad %uint %24 + OpStore %6 %12 + %15 = OpLoad %ulong %4 + %21 = OpIAdd %ulong %15 %ulong_4 + %25 = OpConvertUToPtr %_ptr_Generic_uint %21 + %14 = OpLoad %uint %25 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %17 = OpLoad %uint %7 + %26 = OpConvertUToPtr %_ptr_Generic_uint %16 + OpStore %26 %17 + %18 = OpLoad %ulong %5 + %19 = OpLoad %uint %6 + %23 = OpIAdd %ulong %18 %ulong_4_0 + %27 = OpConvertUToPtr %_ptr_Generic_uint %23 + OpStore %27 %19 + OpReturn + OpFunctionEnd |