diff --git a/README.noformat b/README.noformat
index 8460598..0f2cc85 100644
--- a/README.noformat
+++ b/README.noformat
@@ -41,3 +41,8 @@ Key             _MKN_GPU_THREADED_STREAM_LAUNCHER_WAIT_MS_MAX_
 Type            uint
 Default         100
 Description     Max wait time in milliseconds for polling active jobs for completion when no job is finished.
+
+Key             MKN_CPU_DO_NOT_DEFINE_DIM3
+Type            bool
+Default         false
+Description     if true, skips defining dim3 struct which is usually provided by gpu headers
diff --git a/inc/mkn/gpu.hpp b/inc/mkn/gpu.hpp
index 1f6c813..7a7f978 100644
--- a/inc/mkn/gpu.hpp
+++ b/inc/mkn/gpu.hpp
@@ -35,7 +35,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
 namespace mkn::gpu {
 
-__device__ uint32_t idx() {
+__device__ uint32_t inline idx() {
 #if MKN_GPU_ROCM
   return mkn::gpu::hip::idx();
 
diff --git a/inc/mkn/gpu/cpu.hpp b/inc/mkn/gpu/cpu.hpp
index 15f593f..814a6ed 100644
--- a/inc/mkn/gpu/cpu.hpp
+++ b/inc/mkn/gpu/cpu.hpp
@@ -41,6 +41,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include "mkn/gpu/cli.hpp"
 #include "mkn/gpu/def.hpp"
 
+#include <algorithm>
 #include <cassert>
 #include <cstring>
 
@@ -62,9 +63,11 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #endif
 
 // we need to exclude these for CPU only operations
+#define __shared__
 #define __device__
 #define __host__
 #define __global__
+#define __syncthreads(...)
 
 #if defined(MKN_GPU_FN_PER_NS) && MKN_GPU_FN_PER_NS
 #define MKN_GPU_NS mkn::gpu::cpu
@@ -72,8 +75,11 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #define MKN_GPU_NS mkn::gpu
 #endif  // MKN_GPU_FN_PER_NS
 
-namespace MKN_GPU_NS {
+#if !defined(MKN_CPU_DO_NOT_DEFINE_DIM3)
+#define MKN_CPU_DO_NOT_DEFINE_DIM3 0
+#endif
 
+#if !defined(dim3) and !MKN_CPU_DO_NOT_DEFINE_DIM3
 struct dim3 {
   dim3() {}
   dim3(std::size_t x) : x{x} {}
@@ -83,6 +89,14 @@ struct dim3 {
   std::size_t x = 1, y = 1, z = 1;
 };
 
+dim3 static inline threadIdx, blockIdx;
+
+#endif  // MKN_CPU_DO_NOT_DEFINE_DIM3
+
+//
+
+namespace MKN_GPU_NS {
+
 void inline setLimitMallocHeapSize(std::size_t const& /*bytes*/) {} /*noop*/
 
 void inline setDevice(std::size_t const& /*dev*/) {} /*noop*/
@@ -256,6 +270,16 @@ struct GLauncher : public Launcher {
   std::size_t count;
 };
 
+template <typename Container, typename T>
+void fill(Container& c, size_t const size, T const val) {
+  std::fill(c.begin(), c.begin() + size, val);
+}
+
+template <typename Container, typename T>
+void fill(Container& c, T const val) {
+  fill(c, c.size(), val);
+}
+
 void inline prinfo(std::size_t /*dev*/ = 0) { KOUT(NON) << "Psuedo GPU in use"; }
 
 }  // namespace MKN_GPU_NS
@@ -263,7 +287,7 @@ void inline prinfo(std::size_t /*dev*/ = 0) { KOUT(NON) << "Psuedo GPU in use";
 namespace mkn::gpu::cpu {
 
 template <typename SIZE = std::uint32_t /*max 4294967296*/>
-SIZE idx() {
+SIZE inline idx() {
   return MKN_GPU_NS::detail::idx;
 }
 
diff --git a/inc/mkn/gpu/defines.hpp b/inc/mkn/gpu/defines.hpp
index 533a5ec..3a45d43 100644
--- a/inc/mkn/gpu/defines.hpp
+++ b/inc/mkn/gpu/defines.hpp
@@ -32,7 +32,6 @@
 #endif
 
 #if MKN_GPU_CUDA == 0 && MKN_GPU_ROCM == 0 && !defined(MKN_GPU_CPU)
-#pragma message("mkn.gpu error: No accelerator found, defaulting to CPU IMP")
 #define MKN_GPU_CPU 1
 #endif