Add support for cache_modifier::nc (non-coherent cache) via __ldg

Author: stijnh

docs/guides/cache_modifiers.md

@@ -6,7 +6,7 @@ Kernel Float exposes these hints through the `kf::cache_modifier` enumeration.
 
 ## The `cache_modifier` enum
 
-`kf::cache_modifier` mirrors CUDA's `__ldca`/`__ldcg`/`__ldcs`/`__ldlu`/`__ldcv` load intrinsics and `__stwb`/`__stcg`/`__stcs`/`__stwt` store intrinsics:
+`kf::cache_modifier` mirrors CUDA's `__ldca`/`__ldcg`/`__ldcs`/`__ldlu`/`__ldcv`/`__ldg` load intrinsics and `__stwb`/`__stcg`/`__stcs`/`__stwt` store intrinsics:
 
 | Value                       | Meaning                                                     | Load intrinsic | Store intrinsic |
 |-----------------------------|-------------------------------------------------------------|-----------------|------------------|
@@ -16,11 +16,12 @@ Kernel Float exposes these hints through the `kf::cache_modifier` enumeration.
 | `streaming` (alias `cs`)    | Streaming access, data is likely accessed only once         | `__ldcs`        | `__stcs`         |
 | `uncached` (alias `cv`/`wt`) | Bypass caching entirely and always go to memory             | `__ldcv`        | `__stwt`         |
 | `last_use` (alias `lu`)     | Last use; the cache line will not be reused (loads only)    | `__ldlu`        | -                |
+| `read_only` (alias `nc`)    | Load through the read-only/non-coherent cache (loads only, data must not be written during the kernel) | `__ldg` | -       |
 
-Note that you can use both the short PTX-derived name (`ca`, `cg`, `cs`, `cv`, `wt`, `lu`)
-or a more descriptive alias (`cache_all`, `cache_global`, `streaming`, `uncached`, `last_use`).
+Note that you can use both the short PTX-derived name (`ca`, `cg`, `cs`, `cv`, `wt`, `lu`, `nc`)
+or a more descriptive alias (`cache_all`, `cache_global`, `streaming`, `uncached`, `last_use`, `read_only`).
 
-`last_use` has no store equivalent (there is no `__stlu` intrinsic), so using it with a store operation silently falls back to a plain, unmodified store.
+`last_use` and `read_only` have no store equivalent (there is no `__stlu` or corresponding store for `__ldg`), so using either with a store operation silently falls back to a plain, unmodified store.
 
 ## Using `read_aligned`/`write_aligned`
 
@@ -64,7 +65,7 @@ Cache modifiers are a *hint*, not a guarantee. Kernel Float will attempt to emit
 but will fall back to regular load/store if not possible. There are several cases:
 
 * The code is not being compiled for a CUDA device (e.g., host or HIP). The cache intrinsics are CUDA-specific.
-* The requested modifier is `normal` (or `last_use` for a store, since it has no store equivalent).
+* The requested modifier is `normal` (or `last_use`/`read_only` for a store, since neither has a store equivalent).
 * The data type being accessed has a size that does not match the supported cache intrinsics (1, 2, 4, 8, or 16 bytes). 
 
 In these cases, the cache hint is ignored and a regular load/store instruction is emitted.

include/kernel_float/memory.h

@@ -136,6 +136,9 @@ enum struct cache_modifier {
 
     last_use,  // last use, the cache line will not be re-used afterwards (loads only, not store equivalent).
     lu = last_use,  // Load only: last use, the cache line will not be re-used (`__ldlu`)
+
+    read_only,  // load through the read-only/non-coherent cache (loads only, not store equivalent).
+    nc = read_only  // Load only: non-coherent, read-only cache load (`ld.global.nc`, i.e. `__ldg`).
 };
 
 namespace detail {
@@ -210,6 +213,8 @@ KERNEL_FLOAT_INLINE S cache_load(const S* ptr) {
             result = __ldlu(raw_ptr);
         } else if constexpr (Modifier == cache_modifier::cv) {
             result = __ldcv(raw_ptr);
+        } else if constexpr (Modifier == cache_modifier::nc) {
+            result = __ldg(raw_ptr);
         } else {
             return *ptr;
         }

single_include/kernel_float.h

@@ -16,8 +16,8 @@
 
 //================================================================================
 // this file has been auto-generated, do not modify its contents!
-// date: 2026-07-23 12:03:44.219200
-// git hash: 91cc6e374e07f8c2222c6ded1755cb774b6fb1af
+// date: 2026-08-25 17:05:35.335482
+// git hash: f4e2d65be6e0f28156781eb766bae0fbcbe48b42
 //================================================================================
 
 #ifndef KERNEL_FLOAT_MACROS_H
@@ -2756,6 +2756,9 @@ enum struct cache_modifier {
 
     last_use,  // last use, the cache line will not be re-used afterwards (loads only, not store equivalent).
     lu = last_use,  // Load only: last use, the cache line will not be re-used (`__ldlu`)
+
+    read_only,  // load through the read-only/non-coherent cache (loads only, not store equivalent).
+    nc = read_only  // Load only: non-coherent, read-only cache load (`ld.global.nc`, i.e. `__ldg`).
 };
 
 namespace detail {
@@ -2830,6 +2833,8 @@ KERNEL_FLOAT_INLINE S cache_load(const S* ptr) {
             result = __ldlu(raw_ptr);
         } else if constexpr (Modifier == cache_modifier::cv) {
             result = __ldcv(raw_ptr);
+        } else if constexpr (Modifier == cache_modifier::nc) {
+            result = __ldg(raw_ptr);
         } else {
             return *ptr;
         }

tests/memory.cu

@@ -175,6 +175,10 @@ struct cache_modifier_test {
             auto v = kf::read_aligned<N, kf::cache_modifier::cv>(input.data);
             ASSERT_EQ_ALL(v[I], T(double(I)));
         }
+        {
+            auto v = kf::read_aligned<N, kf::cache_modifier::nc>(input.data);
+            ASSERT_EQ_ALL(v[I], T(double(I)));
+        }
 
         auto v = kf::read_aligned<N>(input.data);
 
@@ -210,6 +214,12 @@ struct cache_modifier_test {
             kf::write_aligned<N, kf::cache_modifier::lu>(output.data, v);
             ASSERT_EQ_ALL(output.data[I], T(double(I)));
         }
+        {
+            // `nc` has no store equivalent; write_aligned should fall back to a plain store.
+            storage_type output = {T(double(I * 0))...};
+            kf::write_aligned<N, kf::cache_modifier::nc>(output.data, v);
+            ASSERT_EQ_ALL(output.data[I], T(double(I)));
+        }
     }
 };
 

tests_codegen/vectorized_load.cu

@@ -87,5 +87,19 @@ __global__ void vector4_load_cv(const float *input, float *output) {
     output[0] = kf::sum(kf::read_aligned<4, kf::cache_modifier::uncached>(input));
 }
 
+// aligned read of 2 elements through the read-only/non-coherent cache, packed
+// into a single scalar 64-bit load (no native .vN form for cache-modified loads)
+// CHECK: ld.global.nc.u64
+__global__ void vector2_load_nc(const float *input, float *output) {
+    output[0] = kf::sum(kf::read_aligned<2, kf::cache_modifier::read_only>(input));
+}
+
+// aligned read of 4 elements through the read-only/non-coherent cache,
+// vectorized as a v2.u64 load
+// CHECK: ld.global.nc.v2.u64
+__global__ void vector4_load_nc(const float *input, float *output) {
+    output[0] = kf::sum(kf::read_aligned<4, kf::cache_modifier::read_only>(input));
+}
+
 
 }