diff --git a/aether/custom_kernels/dropout_kernel.py b/aether/custom_kernels/dropout_kernel.py
index 81ee85a..11692ed 100644
--- a/aether/custom_kernels/dropout_kernel.py
+++ b/aether/custom_kernels/dropout_kernel.py
@@ -1,10 +1,16 @@
+import warnings
+from string import Template
+
+import numpy as np
+
 import aether.config as config
 
 _PHILOX_PREAMBLE = r'''
-__device__ __forceinline__ float philox_uniform(
+__device__ __forceinline__ void philox4x32_10(
         unsigned long long philox_seed,
         unsigned long long philox_offset,
-        long long idx) {
+        long long idx,
+        unsigned int* words) {
 
     unsigned int c0 = (unsigned int)(philox_offset & 0xffffffffULL);
     unsigned int c1 = (unsigned int)(philox_offset >> 32);
@@ -32,87 +38,129 @@
         k0 += 0x9E3779B9u;
         k1 += 0xBB67AE85u;
     }
-    return (c0 >> 8) * (1.0f / 16777216.0f);
+    words[0] = c0; words[1] = c1; words[2] = c2; words[3] = c3;
 }
 '''
 
-_philox_dropout_forward = config.build_kernel(
+# Each thread owns the 8 elements starting at thread_index * 8 and splits one Philox
+# call into their 16-bit draws, low half first. threshold is compared as 32 bits so
+# 65536 (rate 0) keeps every draw. Forward and backward launch this same source, and
+# the draws never see the data, so both directions drop the same elements in any dtype.
+_DROPOUT_TEMPLATE = Template(r'''
+$include
+$philox
+
+extern "C" __global__
+void $kernel_name(
+    const $data_t* __restrict__ src,
+    $data_t* __restrict__ dst,
+    const unsigned long long philox_seed,
+    const unsigned long long philox_offset,
+    const unsigned int threshold,
+    const $work_t inv_keep,
+    const long long size
+) {
+    const long long thread_index = (long long)blockIdx.x * blockDim.x + threadIdx.x;
+    const long long first = thread_index * 8;
+    if (first >= size) return;
+
+    unsigned int words[4];
+    philox4x32_10(philox_seed, philox_offset, thread_index, words);
+
+    for (int j = 0; j < 8 && first + j < size; ++j) {
+        const unsigned int draw = (words[j >> 1] >> ((j & 1) * 16)) & 0xffffu;
+        const long long i = first + j;
+        dst[i] = $store(draw < threshold ? $load(src[i]) * inv_keep : ($work_t)0);
+    }
+}
+''')
+
+_INCLUDES = {
+    "hip": "#include <hip/hip_runtime.h>\n#include <hip/hip_fp16.h>",
+    "cuda": "#include <cuda_fp16.h>",
+}
+
+_DTYPES = {
+    np.dtype(np.float16): dict(data_t='half', work_t='float', load='__half2float', store='__float2half'),
+    np.dtype(np.float32): dict(data_t='float', work_t='float', load='', store=''),
+    np.dtype(np.float64): dict(data_t='double', work_t='double', load='', store=''),
+}
+
+# A dropped channel is exactly 0: multiplying alone would turn NaN or inf into NaN.
+_channel_mask_apply = config.build_kernel(
     lambda: config.cp.ElementwiseKernel(
-        'T x, uint64 philox_seed, uint64 philox_offset, float64 keep_prob',
+        'T x, U mask',
         'T y',
-        '''
-        float u = philox_uniform(philox_seed, philox_offset, i);
-        y = (u < keep_prob) ? (T)(x / keep_prob) : (T)0;
-        ''',
-        '_philox_dropout_forward',
-        preamble=_PHILOX_PREAMBLE,
+        'y = mask != 0 ? (T)(x * mask) : (T)0',
+        'channel_mask_apply_kernel',
     ),
-    name='_philox_dropout_forward'
+    name='_channel_mask_apply',
 )
 
-_philox_dropout_backward = config.build_kernel(
-    lambda: config.cp.ElementwiseKernel(
-        'T dvalues, uint64 philox_seed, uint64 philox_offset, float64 keep_prob',
-        'T dinputs',
-        '''
-        float u = philox_uniform(philox_seed, philox_offset, i);
-        dinputs = (u < keep_prob) ? (T)(dvalues / keep_prob) : (T)0;
-        ''',
-        '_philox_dropout_backward',
-        preamble=_PHILOX_PREAMBLE,
-    ),
-    name='_philox_dropout_backward'
-)
+_dropout_kernel_cache = {}
 
-_philox_spatial_dropout_forward = config.build_kernel(lambda: config.cp.ElementwiseKernel(
-    'T x, uint64 philox_seed, uint64 philox_offset, float64 keep_prob, int64 C, int64 HWC',
-    'T y',
-    '''
-    // Collapse the per-element index into a per-(sample, channel) index.
-    long long c = (i / HWC) * C + i % C;
-    float u = philox_uniform(philox_seed, philox_offset, c);
-    y = (u < keep_prob) ? (T)(x / keep_prob) : (T)0;
-    ''',
-    '_philox_spatial_dropout_forward',
-    preamble=_PHILOX_PREAMBLE,
-    ),
-    name = '_philox_spatial_dropout_forward'
-)
 
-_philox_spatial_dropout_backward = config.build_kernel(lambda: config.cp.ElementwiseKernel(
-    'T dvalues, uint64 philox_seed, uint64 philox_offset, float64 keep_prob, int64 C, int64 HWC',
-    'T dinputs',
-    '''
-    // Compute the same channel mask used in forward pass, as the output is deterministic
-    // passing in our forward inputs results in the same same u value.
-    long long c = (i / HWC) * C + i % C;
-    float u = philox_uniform(philox_seed, philox_offset, c);
-    dinputs = (u < keep_prob) ? (T)(dvalues / keep_prob) : (T)0;
-    ''',
-    '_philox_spatial_dropout_backward',
-    preamble=_PHILOX_PREAMBLE,
-    ),
-    name = '_philox_spatial_dropout_backward'
-)
+def _get_compiled_kernel(variant: str, dtype):
+    """Cached, compiled kernel for one vendor variant and data dtype.
 
-# Helper functions for layer Dropout
-def is_gpu_dropout_available() -> bool:
-    """Checks if CuPy hardware support and kernels are loaded."""
-    return _philox_dropout_forward is not None and _philox_dropout_backward is not None
+    A failed compile is cached as None so a broken build doesn't re-attempt (and
+    re-warn) on every launch.
+    """
+    key = (variant, dtype)
+    if key in _dropout_kernel_cache:
+        return _dropout_kernel_cache[key]
 
-def philox_dropout_forward(inputs, seed: int, offset: int, keep_rate: float):
-    return _philox_dropout_forward(inputs, seed, offset, float(keep_rate)) # type: ignore
+    if dtype not in _DTYPES:
+        raise TypeError(f"[aether] GPU dropout reads float16, float32 or float64, not {dtype}.")
 
-def philox_dropout_backward(dvalues, seed: int, offset: int, keep_rate: float):
-    return _philox_dropout_backward(dvalues, seed, offset, float(keep_rate)) # type: ignore
+    kernel_name = f"philox_dropout_{dtype.name}_{variant}_kernel"
+    source = _DROPOUT_TEMPLATE.substitute(
+        include=_INCLUDES[variant],
+        philox=_PHILOX_PREAMBLE,
+        kernel_name=kernel_name,
+        **_DTYPES[dtype],
+    )
 
-# Helper functions for layer SpatialDropout
-def is_gpu_spatial_dropout_available() -> bool:
-    """Checks if CuPy hardware support and kernels are loaded."""
-    return _philox_spatial_dropout_forward is not None and _philox_spatial_dropout_backward is not None
+    try:
+        kernel = config.cp.RawModule(code=source).get_function(kernel_name)
+    except (config.cp.cuda.compiler.CompileException, config.cp.cuda.driver.CUDADriverError) as e:
+        warnings.warn(f"[aether] dropout_kernel: {kernel_name} compile failed: {e}")
+        kernel = None
 
-def philox_spatial_dropout_forward(inputs, seed: int, offset: int, keep_rate: float, C: int, HWC: int):
-    return _philox_spatial_dropout_forward(inputs, seed, offset, float(keep_rate), C, HWC) # type: ignore
+    _dropout_kernel_cache[key] = kernel
+    return kernel
 
-def philox_spatial_dropout_backward(dvalues, seed: int, offset: int, keep_rate: float, C: int, HWC: int):
-    return _philox_spatial_dropout_backward(dvalues, seed, offset, keep_rate, C, HWC) # type: ignore
+
+def is_gpu_dropout_available() -> bool:
+    """Checks if CuPy hardware support is loaded. Kernels compile at first launch."""
+    return config.HAS_CUPY
+
+
+def apply_channel_mask(x, mask):
+    return _channel_mask_apply(x, mask)
+
+
+def philox_dropout(src, seed: int, offset: int, threshold: int, inv_keep: float):
+    """Returns src with each element kept and scaled by inv_keep, or zeroed.
+
+    The result has src's shape and dtype. Returns None when the kernel for that
+    dtype failed to compile, so the caller can fall back.
+    """
+    variant, threads = config.resolve_gpu_launch_geometry()
+    kernel = _get_compiled_kernel(variant, np.dtype(src.dtype))
+    if kernel is None:
+        return None
+
+    xp = config.cp
+    src = xp.ascontiguousarray(src)
+    dst = xp.empty(src.shape, dtype=src.dtype)
+    size = src.size
+    philox_calls = -(-size // 8)
+    blocks = -(-philox_calls // threads)
+    scale = np.float64(inv_keep) if src.dtype == xp.float64 else np.float32(inv_keep)
+
+    kernel(
+        (blocks,), (threads,),
+        (src, dst, np.uint64(seed), np.uint64(offset), np.uint32(threshold), scale, np.int64(size)),
+    )
+    return dst
diff --git a/aether/layers/dropout.py b/aether/layers/dropout.py
index ff61343..1b2dfc3 100644
--- a/aether/layers/dropout.py
+++ b/aether/layers/dropout.py
@@ -21,6 +21,10 @@ def __init__(self, rate, seed=None):
         self.rate = float(rate)
         self.keep_rate = 1.0 - self.rate
         self._inv_keep = 1.0 / self.keep_rate
+        # The GPU kernel compares 16-bit draws, so its keep rate is threshold / 65536.
+        # Scaling by the inverse of that rate keeps the expectation exact.
+        self._keep_threshold = max(1, round(self.keep_rate * 65536))
+        self._gpu_inv_keep = 65536 / self._keep_threshold
 
         self.seed = seed
         self._seed_key = config.derive_stream_seed(seed, 0)
@@ -57,6 +61,12 @@ def _require_offset(self):
             )
         return self._active_offset
 
+    @staticmethod
+    def _apply_mask(x, mask):
+        """A dropped element is exactly 0: multiplying alone would turn NaN or inf into NaN."""
+        xp = config.get_array_module(x)
+        return xp.where(mask != 0, x, x.dtype.type(0.0)) * mask
+
     def get_config(self):
         return {
             "rate": self.rate,
@@ -73,6 +83,8 @@ class Dropout(_DropoutBase):
     Notes
     -----
     Identity is passed during inference, skipping and dropout calculation.
+    On GPU the rate rounds to the nearest 1/65536 (error at most 7.6e-6), so rates
+    below 7.6e-6 drop nothing. No mask is stored: backward rebuilds it from the seed.
     """
     def _compile_for_device(self, device):
         """Triggered by Model.to(device) to map low-level hardware paths."""
@@ -88,20 +100,32 @@ def _forward_gpu(self, inputs, training):
         if not training:
             self._active_offset = -1
             return inputs
-        
+
         offset = self._clock.value
         self._active_offset = offset
 
-        self.output = gpu_dropout.philox_dropout_forward(
-            inputs, self._seed_key, offset, self.keep_rate
-        )  # type: ignore
+        output = gpu_dropout.philox_dropout(
+            inputs, self._seed_key, offset, self._keep_threshold, self._gpu_inv_keep
+        )
+        if output is None:
+            self.forward = self._forward_fallback
+            self.backward = self._backward_fallback
+            return self._forward_fallback(inputs, training)
+
+        self.output = output
         return self.output
 
     def _backward_gpu(self, dvalues):
 
-        self.dinputs = gpu_dropout.philox_dropout_backward(
-            dvalues, self._seed_key, self._require_offset(), self.keep_rate
-        )  # type: ignore
+        dinputs = gpu_dropout.philox_dropout(
+            dvalues, self._seed_key, self._require_offset(), self._keep_threshold, self._gpu_inv_keep
+        )
+        if dinputs is None:
+            raise RuntimeError(
+                f"[aether] Dropout: failed to compile the GPU kernel for {dvalues.dtype} gradients."
+            )
+
+        self.dinputs = dinputs
         return self.dinputs
 
     def _forward_fallback(self, inputs, training):
@@ -122,12 +146,12 @@ def _forward_fallback(self, inputs, training):
         uniform = gen.random(inputs.shape, dtype=np.float32)
         self.binary_mask = xp.where(uniform < self.keep_rate, keep, drop)
 
-        self.output = inputs * self.binary_mask
+        self.output = self._apply_mask(inputs, self.binary_mask)
         return self.output
 
     def _backward_fallback(self, dvalues):
         self._require_offset()
-        self.dinputs = dvalues * self.binary_mask
+        self.dinputs = self._apply_mask(dvalues, self.binary_mask)
         return self.dinputs
 
 
@@ -142,11 +166,13 @@ class SpatialDropout(_DropoutBase):
     Notes
     -----
     Identity is passed during inference, skipping and dropout calculation.
+    On GPU the rate rounds to the nearest 1/65536 (error at most 7.6e-6), so rates
+    below 7.6e-6 drop nothing. The (S, 1, 1, C) mask stays cached for backward.
     """
 
     def _compile_for_device(self, device):
         """Triggered by Model.to(device) to map low-level hardware paths."""
-        if device == "cupy" and gpu_dropout.is_gpu_spatial_dropout_available():
+        if device == "cupy" and gpu_dropout.is_gpu_dropout_available():
             self.forward = self._forward_gpu
             self.backward = self._backward_gpu
         else:
@@ -157,25 +183,29 @@ def _forward_gpu(self, inputs, training):
 
         if not training:
             self._active_offset = -1
+            self.channel_mask = None
             return inputs
-        
+
         offset = self._clock.value
         self._active_offset = offset
 
-        # Stash the geometry so backward can rebuild the Philox mask.
-        self.C = inputs.shape[-1]
-        self.HWC = inputs.size // inputs.shape[0]
+        xp = config.get_array_module(inputs)
+        ones = xp.ones((inputs.shape[0], 1, 1, inputs.shape[-1]), dtype=xp.float32)
+        mask = gpu_dropout.philox_dropout(
+            ones, self._seed_key, offset, self._keep_threshold, self._gpu_inv_keep
+        )
+        if mask is None:
+            self.forward = self._forward_fallback
+            self.backward = self._backward_fallback
+            return self._forward_fallback(inputs, training)
 
-        self.output = gpu_dropout.philox_spatial_dropout_forward(
-            inputs, self._seed_key, offset, self.keep_rate, self.C, self.HWC
-        )  # type: ignore
+        self.channel_mask = mask.astype(inputs.dtype, copy=False)
+        self.output = gpu_dropout.apply_channel_mask(inputs, self.channel_mask)
         return self.output
 
     def _backward_gpu(self, dvalues):
-
-        self.dinputs = gpu_dropout.philox_spatial_dropout_backward(
-            dvalues, self._seed_key, self._require_offset(), self.keep_rate, self.C, self.HWC
-        )  # type: ignore
+        self._require_offset()
+        self.dinputs = gpu_dropout.apply_channel_mask(dvalues, self.channel_mask)
         return self.dinputs
 
     def _forward_fallback(self, inputs, training):
@@ -198,10 +228,10 @@ def _forward_fallback(self, inputs, training):
         uniform = gen.random((inputs.shape[0], 1, 1, self.C), dtype=np.float32)
         self.channel_mask = xp.where(uniform < self.keep_rate, keep, drop)
 
-        self.output = inputs * self.channel_mask
+        self.output = self._apply_mask(inputs, self.channel_mask)
         return self.output
 
     def _backward_fallback(self, dvalues):
         self._require_offset()
-        self.dinputs = dvalues * self.channel_mask
+        self.dinputs = self._apply_mask(dvalues, self.channel_mask)
         return self.dinputs
\ No newline at end of file
diff --git a/tests/layers/philox_reference.py b/tests/layers/philox_reference.py
new file mode 100644
index 0000000..ad7b1e2
--- /dev/null
+++ b/tests/layers/philox_reference.py
@@ -0,0 +1,29 @@
+"""Naive pure-Python Philox4x32-10, the independent reference for the GPU dropout kernel."""
+MASK32 = 0xffffffff
+
+
+def philox4x32_10(counter, key):
+    c0, c1, c2, c3 = counter
+    k0, k1 = key
+    for _ in range(10):
+        p0 = 0xD2511F53 * c0
+        p1 = 0xCD9E8D57 * c2
+        c0, c1, c2, c3 = (
+            (p1 >> 32) ^ c1 ^ k0,
+            p1 & MASK32,
+            (p0 >> 32) ^ c3 ^ k1,
+            p0 & MASK32,
+        )
+        k0 = (k0 + 0x9E3779B9) & MASK32
+        k1 = (k1 + 0xBB67AE85) & MASK32
+    return c0, c1, c2, c3
+
+
+def kept_elements(seed, offset, size, threshold):
+    """Whether each of `size` flat elements is kept: 8 16-bit draws per Philox call, low half first."""
+    kept = []
+    for thread in range((size + 7) // 8):
+        counter = (offset & MASK32, offset >> 32, thread & MASK32, thread >> 32)
+        for word in philox4x32_10(counter, (seed & MASK32, seed >> 32)):
+            kept += [(word & 0xffff) < threshold, (word >> 16) < threshold]
+    return kept[:size]
diff --git a/tests/layers/test_dropout.py b/tests/layers/test_dropout.py
index 30cfeeb..4abd780 100644
--- a/tests/layers/test_dropout.py
+++ b/tests/layers/test_dropout.py
@@ -1,10 +1,13 @@
 import numpy as np
 
+import aether as ae
 import aether.config as config
 import tests.base_case as base_case
+from aether.custom_kernels import dropout_kernel as gpu_dropout
 from aether.layers.dropout import Dropout
 from aether.layers.linear import Dense
 from aether.model import Model
+from tests.layers.philox_reference import kept_elements, philox4x32_10
 
 
 class TestDropout(base_case.AetherBaseLayerTestCase):
@@ -217,4 +220,202 @@ def test_explicit_seed_wins_over_model_seed(self):
             msg="Same explicit seed=7 at same position must produce same _seed_key regardless of model seed"
         )
 
+    # ---- GPU kernel against the reference ------------------------------------
+
+    def _pinned_layer(self, n, rate):
+        """make_built_layer's seed only reaches build(), so pin the mask stream directly."""
+        layer = self.make_built_layer(Dropout, input_shape=(n,), rate=rate)
+        layer._seed_key = self.FIXED_SEED
+        return layer
+
+    def test_philox_reference_matches_random123_vectors(self):
+        """Published Random123 known answers, so the reference below can be trusted."""
+        self.assertEqual(
+            philox4x32_10((0, 0, 0, 0), (0, 0)),
+            (0x6627e8d5, 0xe169c58d, 0xbc57ac4c, 0x9b00dbd8),
+        )
+        self.assertEqual(
+            philox4x32_10((0x243f6a88, 0x85a308d3, 0x13198a2e, 0x03707344), (0xa4093822, 0x299f31d0)),
+            (0xd16cfe09, 0x94fdcceb, 0x5001e420, 0x24126ea1),
+        )
+
+    def test_gpu_kernel_writes_nothing_past_the_end(self):
+        if not self.uses_gpu_kernel:
+            self.skipTest("Philox kernel only runs on the GPU path")
+
+        variant, threads = config.resolve_gpu_launch_geometry()
+        kernel = gpu_dropout._get_compiled_kernel(variant, np.dtype(np.float32))
+        size = 13
+        src = self.xp.ones(size, dtype=self.xp.float32)
+        dst = self.xp.full(size + 16, -7.0, dtype=self.xp.float32)
+
+        kernel(
+            (1,), (threads,),
+            (src, dst, np.uint64(5), np.uint64(3), np.uint32(32768), np.float32(2.0), np.int64(size)),
+        )
+
+        self.assertEqual(dst[size:].get().tolist(), [-7.0] * 16)
+
+    def test_keep_threshold_and_scale_at_the_rate_limits(self):
+        layer = self.make_built_layer(Dropout, input_shape=(16,), rate=0.1)
+        self.assertEqual(layer._keep_threshold, 58982)
+        self.assertEqual(layer._gpu_inv_keep, 65536 / 58982)
+
+        # rate 0: 65536 only keeps every draw if the kernel compares it as 32 bits.
+        layer = self._pinned_layer(1000, 0.0)
+        self.assertEqual(layer._keep_threshold, 65536)
+        self.assertEqual(layer._gpu_inv_keep, 1.0)
+        inputs = self.xp.linspace(0.1, 5.0, 1000, dtype=self.xp.float32)
+        self.assertTrue(bool(self.xp.all(layer.forward(inputs, training=True) == inputs)))
+
+        # Rounding down to 0 would divide by zero.
+        layer = self._pinned_layer(1000, 0.999999)
+        self.assertEqual(layer._keep_threshold, 1)
+        self.assertEqual(layer._gpu_inv_keep, 65536.0)
+        output = layer.forward(self.xp.ones(1000, dtype=self.xp.float32), training=True)
+        self.assertTrue(bool(self.xp.all(self.xp.isfinite(output))))
+
+    # ---- backward and dtypes -------------------------------------------------
+
+    def test_backward_applies_the_forward_mask_after_the_clock_advances(self):
+        n = 40
+        layer = self._pinned_layer(n, 0.5)
+        kept = layer.forward(self.xp.ones(n, dtype=self.xp.float32), training=True) != 0
+        layer._clock.advance()
+
+        dvalues = self.xp.arange(1, n + 1, dtype=self.xp.float32)
+        dinputs = layer.backward(dvalues)
+
+        # Hand-derived: rate 0.5 scales kept elements by exactly 2.
+        self.assertTrue(bool(self.xp.all(dinputs[kept] == dvalues[kept] * 2.0)))
+        self.assertTrue(bool(self.xp.all(dinputs[~kept] == 0.0)))
+        self.assertTrue(bool(self.xp.any(kept)) and not bool(self.xp.all(kept)))
+
+    def test_output_follows_input_dtype_with_the_same_mask(self):
+        n = 200
+        xp = self.xp
+        masks = {}
+        for dtype in (xp.float16, xp.float32, xp.float64):
+            layer = self._pinned_layer(n, 0.5)
+            output = layer.forward(xp.full(n, 1.5, dtype=dtype), training=True)
+            self.assertEqual(output.dtype, dtype)
+            masks[dtype] = output != 0
+            self.assertTrue(bool(xp.all(output[masks[dtype]] == 3.0)))
+
+        self.assertTrue(bool(xp.all(masks[xp.float16] == masks[xp.float32])))
+        self.assertTrue(bool(xp.all(masks[xp.float64] == masks[xp.float32])))
+
+        # 1 + 2**-40 rounds to 1 in fp32, so a detour through it would lose this.
+        layer = self._pinned_layer(n, 0.5)
+        output = layer.forward(xp.full(n, 1.0 + 2.0**-40, dtype=xp.float64), training=True)
+        self.assertTrue(bool(xp.all(output[masks[xp.float64]] == 2.0 + 2.0**-39)))
+
+    def test_backward_dtype_differing_from_forward_drops_the_same_elements(self):
+        n = 200
+        layer = self._pinned_layer(n, 0.5)
+        kept = layer.forward(self.xp.ones(n, dtype=self.xp.float32), training=True) != 0
+        dinputs = layer.backward(self.xp.ones(n, dtype=self.xp.float16))
+        self.assertTrue(bool(self.xp.all((dinputs != 0) == kept)))
+
+    def test_non_contiguous_input_matches_its_contiguous_copy(self):
+        layer = self._pinned_layer(4, 0.5)
+        inputs = self.xp.arange(1, 241, dtype=self.xp.float32).reshape(60, 4)[::2]
+
+        strided = layer.forward(inputs, training=True)
+        contiguous = layer.forward(inputs.copy(), training=True)
+
+        self.assertEqual(strided.shape, (30, 4))
+        self.assertTrue(bool(self.xp.all(strided == contiguous)))
+        kept = strided != 0
+        self.assertTrue(bool(self.xp.all(strided[kept] == inputs[kept] * 2.0)))
+
+    def test_dropped_elements_are_zero_even_for_nan_and_inf(self):
+        n = 200
+        xp = self.xp
+        layer = self._pinned_layer(n, 0.5)
+        kept = layer.forward(xp.ones(n, dtype=xp.float32), training=True) != 0
+        self.assertTrue(bool(xp.any(kept)) and not bool(xp.all(kept)))
+
+        # Every other element is NaN, the rest +inf, so both land in kept and dropped.
+        bad = xp.full(n, float('inf'), dtype=xp.float32)
+        bad[::2] = float('nan')
+        for result in (layer.forward(bad, training=True), layer.backward(bad)):
+            self.assertTrue(bool(xp.all(result[~kept] == 0.0)))
+            self.assertTrue(bool(xp.all(xp.isnan(result[kept]) == xp.isnan(bad[kept]))))
+            self.assertTrue(bool(xp.all(xp.isinf(result[kept]) == xp.isinf(bad[kept]))))
+
+    # ---- precision: the kernel compiles per dtype, so each one needs its own proof ----
+
+    def _expected_gpu_result(self, values, layer, kept):
+        """Independent of the kernel: scale in fp32, round to the data dtype, zero the dropped."""
+        expected = (values.astype(np.float32) * np.float32(layer._gpu_inv_keep)).astype(values.dtype)
+        expected[~kept] = 0
+        return expected
+
+    def test_gpu_kernel_matches_reference_in_fp16_and_fp32(self):
+        if not self.uses_gpu_kernel:
+            self.skipTest("Philox kernel only runs on the GPU path")
+
+        # A wrong element stride or a missing half conversion would change these bytes.
+        # 8195 crosses a block boundary at both 512 (HIP) and 1024 (CUDA) threads and ends
+        # mid-Philox-call. The large offset exercises the counter's high word. Alternating
+        # dtypes checks the cached variants stay separate.
+        size = 8195
+        offset = 2**32 + 5
+        layer = self._pinned_layer(size, 0.1)
+        layer._clock.value = offset
+        values = np.linspace(-50.0, 50.0, size)
+        kept = np.array(kept_elements(layer._seed_key, offset, size, layer._keep_threshold))
+
+        for dtype in (np.float16, np.float32, np.float16, np.float32):
+            typed = values.astype(dtype)
+            output = layer.forward(self.xp.asarray(typed), training=True)
+            self.assertEqual(output.dtype, dtype)
+            self.assertEqual(output.get().tobytes(), self._expected_gpu_result(typed, layer, kept).tobytes())
+
+    def test_gpu_model_step_runs_the_kernel_in_the_policy_dtype(self):
+        if not self.uses_gpu_kernel:
+            self.skipTest("Philox kernel only runs on the GPU path")
+
+        X = self.xp.asarray(np.random.default_rng(1).normal(size=(64, 32)).astype(np.float32))
+        y = self.xp.zeros(64, dtype=self.xp.int64)
+
+        for precision, dtype in (('float16', np.float16), ('float32', np.float32)):
+            model = Model()
+            model.add(Dense(32, 64))
+            model.add(ae.ReLU())
+            model.add(Dropout(rate=0.2))
+            model.add(Dense(64, 2))
+            model.configure(
+                loss=ae.SoftmaxCategoricalCrossEntropy(),
+                optimizer=ae.Adam(lr=0.001),
+                accuracy=ae.CategoricalAccuracy(),
+            )
+            model.to('cupy')
+            model.manual_seed(seed=42)
+            model.set_precision(compute_dtype=precision)
+            model.finalize(input_shape=(32,))
+            relu, dropout, next_dense = model.layers[1:4]
+
+            out = model.forward(X, training=True)
+            model.loss.calculate(out, y)
+            model.loss.backward(out, y)
+            model.backward(model.loss.dinputs)
+
+            self.assertEqual(dropout.forward.__name__, '_forward_gpu')
+            inputs, dvalues = relu.output.get(), next_dense.dinputs.get()
+            self.assertEqual(inputs.dtype, dtype)
+            self.assertEqual(dvalues.dtype, dtype)
+
+            kept = np.array(
+                kept_elements(dropout._seed_key, dropout._active_offset, inputs.size, dropout._keep_threshold)
+            ).reshape(inputs.shape)
+            self.assertEqual(
+                dropout.output.get().tobytes(), self._expected_gpu_result(inputs, dropout, kept).tobytes()
+            )
+            self.assertEqual(
+                dropout.dinputs.get().tobytes(), self._expected_gpu_result(dvalues, dropout, kept).tobytes()
+            )
+            self.assertTrue(bool((dropout.dinputs.get()[kept] != 0).any()))
+
 base_case.register_test_suites(globals(), TestDropout)
\ No newline at end of file
diff --git a/tests/layers/test_spatial_dropout.py b/tests/layers/test_spatial_dropout.py
index 24bd610..5fb9a07 100644
--- a/tests/layers/test_spatial_dropout.py
+++ b/tests/layers/test_spatial_dropout.py
@@ -1,5 +1,6 @@
 import tests.base_case as base_case
 from aether.layers.dropout import SpatialDropout
+from tests.layers.philox_reference import kept_elements
 
 
 class TestSpatialDropout(base_case.AetherBaseLayerTestCase):
@@ -200,4 +201,70 @@ def test_eval_mode_clears_active_offset(self):
         self.assertEqual(layer._active_offset, -1)
 
 
+    # ---- GPU kernel against the reference ------------------------------------
+
+    def _pinned_layer(self, input_shape, rate):
+        """make_built_layer's seed only reaches build(), so pin the mask stream directly."""
+        layer = self.make_built_layer(SpatialDropout, input_shape=input_shape, rate=rate)
+        layer._seed_key = self.FIXED_SEED
+        return layer
+
+    def test_gpu_mask_matches_reference_per_sample_and_channel(self):
+        if not self.uses_gpu_kernel:
+            self.skipTest("Philox kernel only runs on the GPU path")
+
+        S, C = 9, 13
+        layer = self._pinned_layer((3, 3, C), 0.3)
+        layer._clock.value = 4
+        output = layer.forward(self.xp.ones((S, 3, 3, C), dtype=self.xp.float32), training=True)
+
+        expected = kept_elements(layer._seed_key, 4, S * C, layer._keep_threshold)
+        self.assertEqual((output[:, 0, 0, :] != 0).get().ravel().tolist(), expected)
+
+    def test_backward_applies_the_forward_mask_after_the_clock_advances(self):
+        layer = self._pinned_layer((2, 2, 16), 0.5)
+        shape = (4, 2, 2, 16)
+        kept = layer.forward(self.xp.ones(shape, dtype=self.xp.float32), training=True) != 0
+        layer._clock.advance()
+
+        dvalues = self.xp.arange(1, 257, dtype=self.xp.float32).reshape(shape)
+        dinputs = layer.backward(dvalues)
+
+        # Hand-derived: rate 0.5 scales kept elements by exactly 2.
+        self.assertTrue(bool(self.xp.all(dinputs[kept] == dvalues[kept] * 2.0)))
+        self.assertTrue(bool(self.xp.all(dinputs[~kept] == 0.0)))
+
+    def test_output_follows_input_dtype(self):
+        for dtype in (self.xp.float16, self.xp.float32, self.xp.float64):
+            layer = self._pinned_layer((2, 2, 16), 0.5)
+            output = layer.forward(self.xp.full((4, 2, 2, 16), 1.5, dtype=dtype), training=True)
+            self.assertEqual(output.dtype, dtype)
+            self.assertTrue(bool(self.xp.all(output[output != 0] == 3.0)))
+
+    def test_eval_forward_releases_the_gpu_mask(self):
+        if not self.uses_gpu_kernel:
+            self.skipTest("the cached mask is only released on the GPU path")
+
+        inputs = self.xp.ones((2, 5, 5, 25), dtype=self.xp.float32)
+        self.layer.forward(inputs, training=True)
+        self.assertEqual(self.layer.channel_mask.shape, (2, 1, 1, 25))
+
+        self.layer.forward(inputs, training=False)
+        self.assertIsNone(self.layer.channel_mask)
+
+    def test_dropped_channels_are_zero_even_for_nan_and_inf(self):
+        xp = self.xp
+        shape = (4, 2, 2, 16)
+        layer = self._pinned_layer(shape[1:], 0.5)
+        kept = layer.forward(xp.ones(shape, dtype=xp.float32), training=True) != 0
+        self.assertTrue(bool(xp.any(kept)) and not bool(xp.all(kept)))
+
+        # Every other channel is NaN, the rest +inf, so both land in kept and dropped.
+        bad = xp.full(shape, float('inf'), dtype=xp.float32)
+        bad[..., ::2] = float('nan')
+        for result in (layer.forward(bad, training=True), layer.backward(bad)):
+            self.assertTrue(bool(xp.all(result[~kept] == 0.0)))
+            self.assertTrue(bool(xp.all(xp.isnan(result[kept]) == xp.isnan(bad[kept]))))
+            self.assertTrue(bool(xp.all(xp.isinf(result[kept]) == xp.isinf(bad[kept]))))
+
 base_case.register_test_suites(globals(), TestSpatialDropout)
\ No newline at end of file