diff --git a/src/base/base.h b/src/base/base.h
index 74bc9b56..6625b2a9 100644
--- a/src/base/base.h
+++ b/src/base/base.h
@@ -1,40 +1,28 @@
 #ifndef BASE_H
 #define BASE_H
 
-/* Intrinsic header info:
- * <mmintrin.h"  MMX
- * <xmmintrin.h" SSE
- * <emmintrin.h" SSE2
- * <pmmintrin.h" SSE3
- * <tmmintrin.h" SSSE3
- * <smmintrin.h" SSE4.1
- * <nmmintrin.h" SSE4.2
- * <ammintrin.h" SSE4A
- * <wmmintrin.h" AES
- * <immintrin.h" AVX, AVX2, FMA
-*/
-#include "stddef.h"
-#include "stdint.h"
-#include "stdarg.h"
-#include "intrin.h"
-#include "nmmintrin.h" /* SSE4.2 */
-
-#include "../prof/prof.h"
-
 #include "base_core.h"
-#include "base_intrinsics.h"
-#include "base_atomic.h"
-#include "base_fiber.h"
-#include "base_memory.h"
-#include "base_arena.h"
-#include "base_uid.h"
-#include "base_string.h"
-#include "base_uni.h"
-#include "base_gstat.h"
-#include "base_buddy.h"
-#include "base_math.h"
-#include "base_rand.h"
-#include "base_util.h"
-#include "base_incbin.h"
+
+#if LanguageIsC || LanguageIsCpp
+//- Base cpu includes
+#include "../prof/prof.h"
+# include "base_intrinsics.h"
+# include "base_atomic.h"
+# include "base_fiber.h"
+# include "base_memory.h"
+# include "base_arena.h"
+# include "base_uid.h"
+# include "base_string.h"
+# include "base_uni.h"
+# include "base_gstat.h"
+# include "base_buddy.h"
+# include "base_math.h"
+# include "base_rand.h"
+# include "base_util.h"
+# include "base_incbin.h"
+#elif LanguageIsGpu
+//- Base gpu includes
+# include "base_math_gpu.h"
+#endif
 
 #endif
diff --git a/src/base/base_core.h b/src/base/base_core.h
index 5068c7fa..8a9b618a 100644
--- a/src/base/base_core.h
+++ b/src/base/base_core.h
@@ -58,6 +58,23 @@ extern "C" {
 # error Unknown compiler
 #endif
 
+//- Language
+#if defined(__cplusplus)
+# define LanguageIsCpp 1
+# define LanguageIsC 0
+# define LanguageIsGpu 0
+#elif defined(__STDC_VERSION__)
+# define LanguageIsCpp 0
+# define LanguageIsC 1
+# define LanguageIsGpu 0
+#elif defined(__HLSL_VERSION)
+# define LanguageIsCpp 0
+# define LanguageIsC 0
+# define LanguageIsGpu 1
+#else
+# error Unknown language
+#endif
+
 //- Operating system
 #if defined(_WIN32)
 # define PlatformIsWindows 1
@@ -71,31 +88,30 @@ extern "C" {
 # define PlatformIsWindows 0
 # define PlatformIsMac 0
 # define PlatformIsLinux 1
+#elif LanguageIsGpu
+# define PlatformIsWindows 0
+# define PlatformIsMac 0
+# define PlatformIsLinux 0
 #else
 # error Unknown platform
 #endif
 
-#if defined(__cplusplus)
-# define LanguageIsCpp 1
-# define LanguageIsC 0
-#else
-# define LanguageIsCpp 0
-# define LanguageIsC 1
-#endif
-
 //- Windows NTDDI version
 /* FIXME: Remove this */
 #if 0
-#if PlatformIsWindows
+#if CompilerIsMsvc
 # define NTDDI_WIN11_DT 0x0C0A0000
 # define NTDDI_VERSION 0x0A000000
+# if RtcIsEnabled
+#  define _ALLOW_RTCc_IN_STL 1
+# endif
 #endif
 #endif
 
 ////////////////////////////////
 //~ Debug
 
-//- Compile time assert
+//- Static assert
 #if CompilerIsMsvc || (LanguageIsC && __STDC_VERSION__ < 202311L)
 # if CompilerIsMsvc
 #  define StaticAssert2(cond, line) struct STATIC_ASSERT_____##line {int foo[(cond) ? 1 : -1];}
@@ -108,7 +124,22 @@ extern "C" {
 # define StaticAssert(c) static_assert(c, "")
 #endif
 
+//- Debug assert
+#if RtcIsEnabled
+# if CompilerIsMsvc
+#  define Assert(cond) ((cond) ? 1 : ((*(volatile int *)0) = 0, 0))
+#  define DEBUGBREAK __debugbreak
+# else
+#  define Assert(cond) ((cond) ? 1 : (__builtin_trap(), 0))
+#  define DEBUGBREAK __builtin_debugtrap()
+# endif
+# define DEBUGBREAKABLE { volatile i32 __DEBUGBREAKABLE_VAR = 0; (UNUSED) __DEBUGBREAKABLE_VAR; } (void)0
+#else
+# define Assert(cond) (void)(0)
+#endif
+
 //- Debug alias
+/* TODO: Remove this */
 #if CompilerIsMsvc
 # if DebinfoEnabled
 #  define DebugAlias(var, alias) *(alias) = &(var)
@@ -123,20 +154,6 @@ extern "C" {
 # endif
 #endif
 
-//- Runtime assert
-#if RtcIsEnabled
-# if CompilerIsMsvc
-#  define Assert(cond) ((cond) ? 1 : ((*(volatile int *)0) = 0, 0))
-#  define DEBUGBREAK __debugbreak
-# else
-#  define Assert(cond) ((cond) ? 1 : (__builtin_trap(), 0))
-#  define DEBUGBREAK __builtin_debugtrap()
-# endif
-# define DEBUGBREAKABLE { volatile i32 __DEBUGBREAKABLE_VAR = 0; (UNUSED) __DEBUGBREAKABLE_VAR; } (void)0
-#else
-# define Assert(cond) (void)(0)
-#endif
-
 //- Address sanitization
 #if AsanIsEnabled
 void __asan_poison_memory_region(void const volatile *, size_t);
@@ -148,12 +165,6 @@ void __asan_unpoison_memory_region(void const volatile *add, size_t);
 # define AsanUnpoison(addr, size)
 #endif
 
-//- Allow RTC in STL
-/* Silence Msvc Warning */
-#if RtcIsEnabled && CompilerIsMsvc
-# define _ALLOW_RTCc_IN_STL 1
-#endif
-
 ////////////////////////////////
 //~ Common utility macros
 
@@ -188,8 +199,9 @@ void __asan_unpoison_memory_region(void const volatile *add, size_t);
 
 //- Static
 #define LocalPersist static
-#define internal static
 #define Global static
+/* TODO: Remove this */
+#define internal static
 
 //- Read-only
 #if PlatformIsWindows
@@ -212,6 +224,9 @@ void __asan_unpoison_memory_region(void const volatile *add, size_t);
 #elif defined(__x86_64) || defined(__i386__)
 # define WriteBarrier() __asm__ volatile("" ::: "memory")
 # define ReadBarrier() __asm__ volatile("" ::: "memory")
+#elif LanguageIsGpu
+# define WriteBarrier()
+# define ReadBarrier()
 #else
 # error Memory barriers not implemented
 #endif
@@ -290,7 +305,7 @@ void __asan_unpoison_memory_region(void const volatile *add, size_t);
 #define countof(a) (sizeof(a) / sizeof((a)[0]))
 
 //- IsArray
-#define IsIndexable(a) (sizeof(a[0]))
+#define IsIndexable(a) (sizeof(a[0]) != 0)
 #define IsArray(a) (IsIndexable(a) && (((void *)&a) == ((void *)a)))
 
 ////////////////////////////////
@@ -299,8 +314,10 @@ void __asan_unpoison_memory_region(void const volatile *add, size_t);
 //- Pack
 #if CompilerIsMsvc
 # define Packed(s) __pragma(pack(push, 1)) s __pragma(pack(pop))
-#else
+#elif CompilerIsClang
 # define Packed(s) s __attribute((__packed__))
+#elif LanguageIsGpu
+# define Packed(s) s
 #endif
 
 //- alignas
@@ -337,6 +354,26 @@ void __asan_unpoison_memory_region(void const volatile *add, size_t);
 #define ColorOrange    Rgb32(0xFF, 0xA5, 0x00)
 #define ColorPurple    Rgb32(0xFF, 0x00, 0XFF)
 
+////////////////////////////////
+//~ Intrinsic headers
+
+#if !LanguageIsGpu
+/* Intrinsic header info:
+ * <mmintrin.h"  MMX
+ * <xmmintrin.h" SSE
+ * <emmintrin.h" SSE2
+ * <pmmintrin.h" SSE3
+ * <tmmintrin.h" SSSE3
+ * <smmintrin.h" SSE4.1
+ * <nmmintrin.h" SSE4.2
+ * <ammintrin.h" SSE4A
+ * <wmmintrin.h" AES
+ * <immintrin.h" AVX, AVX2, FMA
+*/
+#include "intrin.h"
+#include "nmmintrin.h" /* SSE4.2 */
+#endif
+
 ////////////////////////////////
 //~ Struct helper macros
 
@@ -346,6 +383,9 @@ void __asan_unpoison_memory_region(void const volatile *add, size_t);
 ////////////////////////////////
 //~ Scalar types
 
+#if !LanguageIsGpu
+//- Cpu scalar types
+#include "stdint.h"
 typedef int8_t      i8;
 typedef int16_t     i16;
 typedef int32_t     i32;
@@ -357,8 +397,16 @@ typedef uint64_t    u64;
 typedef float       f32;
 typedef double      f64;
 typedef i8          b8;
-typedef i32         b32;
+typedef u32         b32;
+#else
+//- Gpu scalar types
+typedef int         i32;
+typedef uint        u32;
+typedef float       f32;
+typedef uint        b32;
+#endif
 
+//- Min / max constants
 #define U8Max          (0xFF)
 #define U16Max         (0xFFFF)
 #define U32Max         (0xFFFFFFFF)
@@ -374,6 +422,8 @@ typedef i32         b32;
 #define I32Min         ((i32)0x80000000)
 #define I64Min         ((i64)0x8000000000000000LL)
 
+//- Float infinity / nan constants
+#if !LanguageIsGpu
 Global const u32 _f32_infinity_u32  = 0x7f800000;
 Global const f32 *_f32_infinity     = (f32 *)&_f32_infinity_u32;
 #define F32Infinity                   (*_f32_infinity)
@@ -392,6 +442,7 @@ Global const f64 *_f64_nan          = (f64 *)&_f64_nan_u64;
 
 #define IsF32Nan(x) (x != x)
 #define IsF64Nan(x) (x != x)
+#endif
 
 ////////////////////////////////
 //~ Config
diff --git a/src/base/base_math.h b/src/base/base_math.h
index 7d1dcf1b..c0a59438 100644
--- a/src/base/base_math.h
+++ b/src/base/base_math.h
@@ -2,6 +2,7 @@
 
 #define Pi              ((f32)3.14159265358979323846)
 #define Tau             ((f32)6.28318530717958647693)
+#define GoldenRatio     ((f32)1.61803398874989484820)
 
 ////////////////////////////////
 //~ Floating point vector2 types
diff --git a/src/base/base_math_gpu.h b/src/base/base_math_gpu.h
new file mode 100644
index 00000000..288d465b
--- /dev/null
+++ b/src/base/base_math_gpu.h
@@ -0,0 +1,19 @@
+#define Pi              ((f32)3.14159265358979323846)
+#define Tau             ((f32)6.28318530717958647693)
+#define GoldenRatio     ((f32)1.61803398874989484820)
+
+#if LanguageIsC || LanguageIsCpp || !LanguageIsGpu
+# error AAA
+#endif
+
+typedef float2 Vec2;
+typedef float3 Vec3;
+typedef float4 Vec4;
+typedef int2 Vec2I32;
+typedef int3 Vec3I32;
+typedef float2x3 Xform;
+typedef float4 Rect;
+typedef float2 ClipRect;
+typedef float2 Aabb;
+typedef float4 Quad;
+typedef float4x4 Mat4x4;
diff --git a/src/gpu/gpu_core_dx12.c b/src/gpu/gpu_core_dx12.c
index c0d8683f..f018ef8a 100644
--- a/src/gpu/gpu_core_dx12.c
+++ b/src/gpu/gpu_core_dx12.c
@@ -2879,9 +2879,9 @@ GPU_Resource *gp_run_render(GPU_RenderSig *gp_render_sig, GPU_RenderParams param
 
             /* Set sig */
             struct k_material_sig sig = ZI;
-            sig.projection = K_Float4x4FromMat4x4(world_to_render_vp_matrix);
-            sig.instances_urid = K_UintFromU32(material_instance_buffer->resource->srv_descriptor->index);
-            sig.grids_urid = K_UintFromU32(grid_buffer->resource->srv_descriptor->index);
+            sig.projection = world_to_render_vp_matrix;
+            sig.instances_urid = material_instance_buffer->resource->srv_descriptor->index;
+            sig.grids_urid = grid_buffer->resource->srv_descriptor->index;
             command_list_set_sig(cl, &sig, sizeof(sig));
 
             /* Draw */
diff --git a/src/kernel/kernel.h b/src/kernel/kernel.h
index 1cf89d2c..6bab975f 100644
--- a/src/kernel/kernel.h
+++ b/src/kernel/kernel.h
@@ -1,6 +1,8 @@
 #ifndef KERNEL_H
 #define KERNEL_H
 
+#include "../base/base.h"
+
 #include "kernel_core.h"
 
 #endif
diff --git a/src/kernel/kernel_core.h b/src/kernel/kernel_core.h
index ed22724c..2180d1e8 100644
--- a/src/kernel/kernel_core.h
+++ b/src/kernel/kernel_core.h
@@ -1,5 +1,5 @@
 /* Determine if file was included from C or from HLSL */
-#if defined(LanguageIsC) || defined(LanguageIsCpp)
+#if !LanguageIsGpu
 # define K_IS_CPU 1
 #else
 # define K_IS_CPU 0
@@ -96,11 +96,6 @@ Inline struct K_float2x3 K_Float2x3FromXform(Xform v)
 
 #define DECLS(t, n) t n : n
 
-#define Tau     6.28318530718
-#define Pi      3.14159265359
-#define GOLDEN  1.61803398875
-
-
 #define resource_from_urid(urid) ResourceDescriptorHeap[urid]
 #define resource_from_nurid(nurid) ResourceDescriptorHeap[NonUniformResourceIndex(nurid)]
 
@@ -166,6 +161,7 @@ SamplerState s_point_clamp : register(s0);
  * Material shader structs
  * ========================== */
 
+#if 0
 K_STRUCT(k_material_sig {
     /* ----------------------------------------------------- */
     K_DECL(float4x4, projection);              /* 16 consts */
@@ -177,6 +173,20 @@ K_STRUCT(k_material_sig {
     /* ----------------------------------------------------- */
 });
 K_ASSERT_ROOT_CONST(struct k_material_sig, 20);
+#else
+Struct(k_material_sig)
+{
+    /* ----------------------------------------------------- */
+    Mat4x4 projection;                          /* 16 consts */
+    /* ----------------------------------------------------- */
+    u32 instances_urid;                         /* 01 consts */
+    u32 grids_urid;                             /* 01 consts */
+    u32 _pad0;                                  /* 01 consts (padding) */
+    u32 _pad1;                                  /* 01 consts (padding) */
+    /* ----------------------------------------------------- */
+};
+K_ASSERT_ROOT_CONST(struct k_material_sig, 20);
+#endif
 
 K_STRUCT(k_material_instance {
     K_DECL(uint, tex_nurid);
diff --git a/src/sprite/sprite_core.c b/src/sprite/sprite_core.c
index f8b92a92..5a9af689 100644
--- a/src/sprite/sprite_core.c
+++ b/src/sprite/sprite_core.c
@@ -233,7 +233,9 @@ S_StartupReceipt sprite_startup(void)
     P_Run(1, sprite_evictor_job, 0, P_Pool_Background, P_Priority_Low, &G.shutdown_counter);
 
     P_OnExit(&sprite_shutdown);
+#if RESOURCE_RELOADING
     W_RegisterCallback(&sprite_watch_callback);
+#endif
 
     return (S_StartupReceipt) { 0 };
 }