aboutsummaryrefslogtreecommitdiffhomepage
path: root/zluda/tests/kernel_suld.ptx
diff options
context:
space:
mode:
authorAndrzej Janik <[email protected]>2021-02-27 20:55:19 +0100
committerAndrzej Janik <[email protected]>2024-02-11 20:45:51 +0100
commit1b9ba2b2333746c5e2b05a2bf24fa6ec3828dcdf (patch)
tree0b77ca4a41d4f232bd181e2bddc886475c608784 /zluda/tests/kernel_suld.ptx
parent60d2124a16a7a2a1a6be3707247afe82892a4163 (diff)
downloadZLUDA-1b9ba2b2333746c5e2b05a2bf24fa6ec3828dcdf.tar.gz
ZLUDA-1b9ba2b2333746c5e2b05a2bf24fa6ec3828dcdf.zip
Nobody expects the Red Teamv3
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
Diffstat (limited to 'zluda/tests/kernel_suld.ptx')
-rw-r--r--zluda/tests/kernel_suld.ptx36
1 files changed, 36 insertions, 0 deletions
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;
+}