diff options
| author | Andrzej Janik <[email protected]> | 2021-03-03 00:59:47 +0100 |
|---|---|---|
| committer | Andrzej Janik <[email protected]> | 2021-03-03 00:59:47 +0100 |
| commit | cdac38d572e2cd86036cbd85f753214b8e1a5172 (patch) | |
| tree | e7cd1536e7a5e90b548f4159a7b196fb8e9326ce /ptx/src/test/spirv_run | |
| parent | 648035a01a84cd87b7f917b277e4e2faad7bb731 (diff) | |
| download | ZLUDA-cdac38d572e2cd86036cbd85f753214b8e1a5172.tar.gz ZLUDA-cdac38d572e2cd86036cbd85f753214b8e1a5172.zip | |
Support kernel tuning directives
Diffstat (limited to 'ptx/src/test/spirv_run')
| -rw-r--r-- | ptx/src/test/spirv_run/add_tuning.ptx | 24 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/add_tuning.spvtxt | 48 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/mod.rs | 1 |
3 files changed, 73 insertions, 0 deletions
diff --git a/ptx/src/test/spirv_run/add_tuning.ptx b/ptx/src/test/spirv_run/add_tuning.ptx new file mode 100644 index 0000000..2a5dcf8 --- /dev/null +++ b/ptx/src/test/spirv_run/add_tuning.ptx @@ -0,0 +1,24 @@ +.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry add_tuning(
+ .param .u64 input,
+ .param .u64 output
+)
+.maxntid 256, 1, 1
+.minnctapersm 4
+{
+ .reg .u64 in_addr;
+ .reg .u64 out_addr;
+ .reg .u64 temp;
+ .reg .u64 temp2;
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+
+ ld.u64 temp, [in_addr];
+ add.u64 temp2, temp, 1;
+ st.u64 [out_addr], temp2;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/add_tuning.spvtxt b/ptx/src/test/spirv_run/add_tuning.spvtxt new file mode 100644 index 0000000..173e0d4 --- /dev/null +++ b/ptx/src/test/spirv_run/add_tuning.spvtxt @@ -0,0 +1,48 @@ + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %23 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "add_tuning" + OpExecutionMode %1 MaxWorkgroupSizeINTEL 256 1 1 + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %26 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong +%_ptr_Generic_ulong = OpTypePointer Generic %ulong + %ulong_1 = OpConstant %ulong 1 + %1 = OpFunction %void None %26 + %8 = OpFunctionParameter %ulong + %9 = OpFunctionParameter %ulong + %21 = 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 + %10 = OpLoad %ulong %2 Aligned 8 + OpStore %4 %10 + %11 = OpLoad %ulong %3 Aligned 8 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %19 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %19 Aligned 8 + OpStore %6 %12 + %15 = OpLoad %ulong %6 + %14 = OpIAdd %ulong %15 %ulong_1 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %17 = OpLoad %ulong %7 + %20 = OpConvertUToPtr %_ptr_Generic_ulong %16 + OpStore %20 %17 Aligned 8 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index 91e6113..4178e2f 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/ptx/src/test/spirv_run/mod.rs @@ -152,6 +152,7 @@ test_ptx!(shared_ptr_take_address, [97815231u64], [97815231u64]); // For now, we just make sure that it builds and links
test_ptx!(assertfail, [716523871u64], [716523872u64]);
test_ptx!(cvt_s64_s32, [-1i32], [-1i64]);
+test_ptx!(add_tuning, [2u64], [3u64]);
struct DisplayError<T: Debug> {
err: T,
|