From 9c606162760cca1e49b333acb3b7ec18e2f7b003 Mon Sep 17 00:00:00 2001
From: Erik Schultheis <erik.schultheis@aalto.fi>
Date: Thu, 26 Jun 2025 17:52:20 +0200
Subject: [PATCH] fix out-of-bounds access in rmsnorm kernel

---
 llmc/rmsnorm.cuh | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/llmc/rmsnorm.cuh b/llmc/rmsnorm.cuh
index 8f20e98..5be05bb 100644
--- a/llmc/rmsnorm.cuh
+++ b/llmc/rmsnorm.cuh
@@ -90,7 +90,7 @@ __global__ void fused_residual_rmsnorm_forward_kernel5(floatX* residual, floatX*
     __syncthreads();
 
     int idx = blockIdx.x * blockDim.y + threadIdx.y;
-    if(idx > N) return;
+    if(idx >= N) return;
 
     // adjust pointers to current token
     residual += C * idx;