updates

.github/workflows/build.yml

@@ -18,6 +18,6 @@ jobs:
 
     - name: test
       run: |
-        curl -Lo mkn https://github.com/mkn/mkn/releases/download/latest/mkn_nix
+        curl -fL --retry 3 --retry-delay 2 -o mkn https://github.com/mkn/mkn/releases/download/latest/mkn_nix
         chmod +x mkn
         KLOG=3 ./mkn clean build run -dtKOgp cpu -a "-std=c++20" test -W 9

.gitignore

@@ -1,4 +1,4 @@
 bin
 .mkn
 *.cui
-.clangd
+*.clangd

inc/mkn/gpu.hpp

@@ -31,7 +31,19 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #ifndef _MKN_GPU_HPP_
 #define _MKN_GPU_HPP_
 
-#include "mkn/gpu/defines.hpp"
+#include "mkn/gpu/def.hpp"
+
+#if MKN_GPU_ROCM
+#include "mkn/gpu/rocm.hpp"
+#endif
+
+#if MKN_GPU_CUDA
+#include "mkn/gpu/cuda.hpp"
+#endif
+
+#if MKN_GPU_CPU
+#include "mkn/gpu/cpu.hpp"
+#endif
 
 namespace mkn::gpu {
 

inc/mkn/gpu/any/cls.hpp

@@ -0,0 +1,24 @@
+#ifndef _MKN_GPU_ANY_CLS_HPP_
+#define _MKN_GPU_ANY_CLS_HPP_
+
+#include "mkn/kul/env.hpp"
+#include "mkn/kul/string.hpp"
+
+namespace mkn::gpu {
+
+template <typename Device>
+struct Cli {
+  constexpr static inline char const* MKN_GPU_BX_THREADS = "MKN_GPU_BX_THREADS";
+
+  auto bx_threads() const {
+    if (kul::env::EXISTS(MKN_GPU_BX_THREADS))
+      return kul::String::INT32(kul::env::GET(MKN_GPU_BX_THREADS));
+    return dev.maxThreadsPerBlock;
+  }
+
+  Device const& dev;
+};
+
+} /* namespace mkn::gpu */
+
+#endif /*_MKN_GPU_ANY_CLS_HPP_*/

inc/mkn/gpu/any/def.hpp

@@ -0,0 +1,4 @@
+#ifndef _MKN_GPU_ANY_DEF_HPP_
+#define _MKN_GPU_ANY_DEF_HPP_
+
+#endif /*_MKN_GPU_ANY_DEF_HPP_*/

inc/mkn/gpu/any/inc/devfunc.ipp

@@ -0,0 +1,9 @@
+
+template <bool sync = true, typename T>
+void __device__ fill_warp_size(T* const t, std::size_t const size, T const val) {
+  std::size_t chunk = 0;
+  auto const each = size / warpSize;
+  for (; chunk < each; ++chunk) t[chunk * warpSize + threadIdx.x] = val;
+  if (threadIdx.x < size - (warpSize * each)) t[chunk * warpSize + threadIdx.x] = val;
+  if constexpr (sync) __syncthreads();
+}

inc/mkn/gpu/device.hpp → inc/mkn/gpu/any/inc/device.ipp

@@ -59,12 +59,12 @@ struct DeviceMem {
   }
 
   void send(T const* t, std::size_t _size = 1, std::size_t start = 0) {
-    MKN_GPU_NS::send(p, t, _size, start);
+    MKN_GPU_NS::send(p + start, t, _size);
   }
 
   template <typename C, std::enable_if_t<mkn::kul::is_span_like_v<C>, bool> = 0>
   void send(C const& c, std::size_t start = 0) {
-    send(c.data(), c.size(), start);
+    send(c.data() + start, c.size());
   }
 
   void fill_n(T t, std::size_t _size, std::size_t start = 0) {
@@ -128,14 +128,14 @@ struct AsioDeviceMem {
 
   void send(Stream& stream, T* t, std::size_t _size = 1, std::size_t start = 0) {
     assert(p != nullptr);
-    MKN_GPU_NS::send_async(p, t, stream, _size, start);
+    MKN_GPU_NS::send_async(p + start, t + start, stream(), _size);
   }
 
   template <typename Span>
   void take(Stream& stream, Span& span, std::size_t start) {
     assert(p != nullptr);
     assert(span.size() + start <= s);
-    MKN_GPU_NS::take_async(p, span, stream, start);
+    MKN_GPU_NS::take_async(p + start, span.data(), stream(), span.size());
   }
 
   auto& size() const { return s; }

inc/mkn/gpu/cpu.hpp

@@ -193,34 +193,32 @@ void send(void* p, void* t, Size size = 1) {
 }
 
 template <typename T, typename Size>
-void send(T* p, T const* t, Size size = 1, Size start = 0) {
+void send(T* p, T const* t, Size size = 1) {
   KLOG(TRC);
-  MKN_GPU_ASSERT(std::memcpy(p + start, t, size * sizeof(T)));
+  MKN_GPU_ASSERT(std::memcpy(p, t, size * sizeof(T)));
 }
-
 template <typename T, typename Size>
-void take(T* p, T* t, Size size = 1, Size start = 0) {
+void send_async(T* p, T const* t, auto& /*stream*/, Size size = 1) {
   KLOG(TRC);
-  MKN_GPU_ASSERT(std::memcpy(t, p + start, size * sizeof(T)));
+  send(p, t, size);
 }
 
 template <typename T, typename Size>
-void send_async(T* p, T const* t, Stream& /*stream*/, Size size = 1, Size start = 0) {
+void take(T const* p, T* t, Size size = 1) {
   KLOG(TRC);
-  send(p, t, size, start);
+  MKN_GPU_ASSERT(std::memcpy(t, p, size * sizeof(T)));
 }
 
-template <typename T, typename Span>
-void take_async(T* p, Span& span, Stream& /*stream*/, std::size_t start) {
-  static_assert(mkn::kul::is_span_like_v<Span>);
+template <typename T, typename Size>
+void take_async(T const* p, T* t, auto& /*stream*/, Size size = 1) {
   KLOG(TRC);
-  take(p, span.data(), span.size(), start);
+  take(p, t, size);
 }
 
 void inline sync() {}
 
-#include "mkn/gpu/alloc.hpp"
-#include "mkn/gpu/device.hpp"
+#include "mkn/gpu/any/inc/alloc.ipp"
+#include "mkn/gpu/any/inc/device.ipp"
 
 namespace detail {
 static thread_local std::size_t idx = 0;
@@ -281,8 +279,8 @@ void fill(Container& c, T const val) {
 }
 
 template <typename T>
-void zero(T* const t, std::size_t const size) {
-  std::fill(t, t + size, 0);
+void fill_warp_size(T* const t, std::size_t const size, T const val) {
+  std::fill(t, t + size, val);
 }
 
 void inline prinfo(std::size_t /*dev*/ = 0) { KOUT(NON) << "Psuedo GPU in use"; }
@@ -310,7 +308,7 @@ static void global_d_kernel(F& f, Args... args) {
   f(args...);
 }
 
-#include "launchers.hpp"
+#include "mkn/gpu/any/inc/launchers.ipp"
 
 } /* namespace MKN_GPU_NS */
 

inc/mkn/gpu/cpu/api.hpp

@@ -0,0 +1,4 @@
+#ifndef _MKN_GPU_CPU_API_HPP_
+#define _MKN_GPU_CPU_API_HPP_
+
+#endif /*_MKN_GPU_CPU_API_HPP_*/

inc/mkn/gpu/cpu/cls.hpp

@@ -0,0 +1,4 @@
+#ifndef _MKN_GPU_CPU_CLS_HPP_
+#define _MKN_GPU_CPU_CLS_HPP_
+
+#endif /*_MKN_GPU_CPU_CLS_HPP_*/

inc/mkn/gpu/cpu/def.hpp

@@ -0,0 +1,4 @@
+#ifndef _MKN_GPU_CPU_DEF_HPP_
+#define _MKN_GPU_CPU_DEF_HPP_
+
+#endif /*_MKN_GPU_CPU_DEF_HPP_*/