From 87fa5af70a488d561f751df546c2de2a7a7956cc Mon Sep 17 00:00:00 2001
From: George Hotz <george@comma.ai>
Date: Tue, 9 May 2023 13:11:56 -0700
Subject: [PATCH] ptx example

---
 extra/ptx/test.py | 23 +++++++++++++++++++++++
 1 file changed, 23 insertions(+)
 create mode 100644 extra/ptx/test.py

diff --git a/extra/ptx/test.py b/extra/ptx/test.py
new file mode 100644
index 0000000000..f30348b8c4
--- /dev/null
+++ b/extra/ptx/test.py
@@ -0,0 +1,23 @@
+#!/usr/bin/env python3
+import numpy as np
+from tinygrad.runtime.ops_cuda import CUDAProgram, RawCUDABuffer
+
+if __name__ == "__main__":
+  test = RawCUDABuffer.fromCPU(np.zeros(10, np.float32))
+  prg = CUDAProgram("test", """
+  .version 7.8
+  .target sm_86
+  .address_size 64
+  .visible .entry test(.param .u64 x) {
+    .reg .b32       %r<2>;
+    .reg .b64       %rd<3>;
+
+    ld.param.u64    %rd1, [x];
+    cvta.to.global.u64      %rd2, %rd1;
+    mov.u32         %r1, 0x40000000; // 2.0 in float
+    st.global.u32   [%rd2], %r1;
+    ret;
+  }""", binary=True)
+  prg([1], [1], test)
+  print(test.toCPU())
+