From 1b9ba2b2333746c5e2b05a2bf24fa6ec3828dcdf Mon Sep 17 00:00:00 2001
From: Andrzej Janik <vosen@vosen.pl>
Date: Sat, 27 Feb 2021 20:55:19 +0100
Subject: Nobody expects the Red Team

Too many changes to list, but broadly:
* Remove Intel GPU support from the compiler
* Add AMD GPU support to the compiler
* Remove Intel GPU host code
* Add AMD GPU host code
* More device instructions. From 40 to 68
* More host functions. From 48 to 184
* Add proof of concept implementation of OptiX framework
* Add minimal support of cuDNN, cuBLAS, cuSPARSE, cuFFT, NCCL, NVML
* Improve ZLUDA launcher for Windows
---
 zluda/tests/kernel_suld.ptx | 36 ++++++++++++++++++++++++++++++++++++
 1 file changed, 36 insertions(+)
 create mode 100644 zluda/tests/kernel_suld.ptx

(limited to 'zluda/tests/kernel_suld.ptx')

diff --git a/zluda/tests/kernel_suld.ptx b/zluda/tests/kernel_suld.ptx
new file mode 100644
index 0000000..4e9b5b1
--- /dev/null
+++ b/zluda/tests/kernel_suld.ptx
@@ -0,0 +1,36 @@
+.version 6.5
+.target sm_30
+.address_size 64
+
+.global .surfref image;
+
+.visible .entry suld(
+    .param .b64 output,
+    .param .b32 input_x,
+    .param .b32 input_y,
+    .param .b32 input_z,
+    .param .b64 image_bindless_param
+)
+{
+    .reg .b32           coord_x;
+    .reg .b32           coord_y;
+    .reg .b32           coord_z;
+    .reg .b32           coord_depth;
+    .reg .u64 	        out_addr;
+    .reg .u64 	        image_bindless;
+
+    ld.param.b32        coord_x, [input_x];
+    ld.param.b32        coord_y, [input_y];
+    ld.param.b32        coord_z, [input_z];
+    ld.param.u64 	    out_addr, [output];
+    ld.param.u64 	    image_bindless, [image_bindless_param];
+    mov.b32             coord_depth, coord_z;
+
+    #REG_VALUES#
+
+    suld.b.#GEOMETRY##FORMAT#.trap #VALUES#, [#IMAGE_SRC#, #COORDINATES#];
+
+    st#FORMAT#          [out_addr], #VALUES#;
+
+    ret;
+}
-- 
cgit v1.2.3