From 1b9ba2b2333746c5e2b05a2bf24fa6ec3828dcdf Mon Sep 17 00:00:00 2001 From: Andrzej Janik 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_texobj_2d.ptx | 34 ++++++++++++++++++++++++++++++++++ 1 file changed, 34 insertions(+) create mode 100644 zluda/tests/kernel_texobj_2d.ptx (limited to 'zluda/tests/kernel_texobj_2d.ptx') diff --git a/zluda/tests/kernel_texobj_2d.ptx b/zluda/tests/kernel_texobj_2d.ptx new file mode 100644 index 0000000..6b1d7db --- /dev/null +++ b/zluda/tests/kernel_texobj_2d.ptx @@ -0,0 +1,34 @@ +.version 6.5 +.target sm_30 +.address_size 64 + +.visible .entry texobj( + .param .f32 input_x, + .param .f32 input_y, + .param .u64 image_param, + .param .u64 output +) +{ + .reg .u64 out_addr; + .reg .u64 temp; + .reg .u64 temp2; + .reg .u64 image; + .reg .f32 x; + .reg .f32 y; + .reg .s32 r; + .reg .s32 g; + .reg .s32 b; + .reg .s32 a; + + ld.param.f32 x, [input_x]; + ld.param.f32 y, [input_y]; + ld.param.u64 image, [image_param]; + ld.param.u64 out_addr, [output]; + + tex.2d.v4.s32.f32 {r, g, b, a}, [image, {x, y}]; + st.b32 [out_addr], a; + st.b32 [out_addr+4], b; + st.b32 [out_addr+8], g; + st.b32 [out_addr+12], r; + ret; +} -- cgit v1.2.3