```json
{
  "abi": "nodus-cuda-rmsnorm/1",
  "block_x": 256,
  "kernel": "nodus_rmsnorm_bf16",
  "shared_bytes": 0,
  "repaired_source": "__global__ void nodus_rmsnorm_bf16(const __nv_bfloat16* x, const __nv_bfloat16* w, __nv_bfloat16* out, int T, int H, float eps) {\n  if (blockIdx.x >= static_cast<unsigned int>(T)) return;\n  const unsigned int tid = threadIdx.x;\n  const unsigned int lane = tid & 31u;\n  const unsigned int warp = tid >> 5;\n  const unsigned int num_warps = blockDim.x >> 5;\n  const unsigned long long row_offset = static_cast<unsigned long long>(blockIdx.x) * static_cast<unsigned long long>(H);\n  __shared__ float warp_sums[32];\n\n  float sum_sq = 0.0f;\n  for (unsigned int col = tid; col < static_cast<unsigned int>(H); col += blockDim.x) {\n    const float value = __bfloat162float(x[row_offset + col]);\n    sum_sq += value * value;\n  }\n\n  for (int offset = 16; offset > 0; offset /= 2) {\n    sum_sq += __shfl_xor_sync(0xffffffffu, sum_sq, offset, 32);\n  }\n  if (lane == 0) warp_sums[warp] = sum_sq;\n  __syncthreads();\n  if (warp == 0) {\n    sum_sq = lane < num_warps ? warp_sums[lane] : 0.0f;\n    for (int offset = 16; offset > 0; offset /= 2) {\n      sum_sq += __shfl_xor_sync(0xffffffffu, sum_sq, offset, 32);\n    }\n    if (lane == 0) warp_sums[0] = sum_sq;\n  }\n\n  const float rms_rcp = rsqrtf(warp_sums[0] / static_cast<float>(blockDim.x) + eps);\n  for (unsigned int col = tid; col < static_cast<unsigned int>(H); col += blockDim.x) {\n    const float value = __bfloat162float(x[row_offset + col]);\n    const float weight = __bfloat162float(w[col]);\n    out[row_offset + col] = __float2bfloat16_rn(value * rms_rcp * weight);\n  }\n}\n"
}
```