{
  "launch": {
    "abi": "nodus-cuda-rmsnorm/1",
    "block_x": 256,
    "kernel": "nodus_rmsnorm_bf16",
    "shared_bytes": 0
  },
  "protocol": "nodus-cuda-cpp-submission/1",
  "source": "/*\n * Copyright (c) 2024 by FlashInfer team.\n *\n * Licensed under the Apache License, Version 2.0 (the \"License\");\n * you may not use this file except in compliance with the License.\n * You may obtain a copy of the License at\n *\n *   http://www.apache.org/licenses/LICENSE-2.0\n *\n * Unless required by applicable law or agreed to in writing, software\n * distributed under the License is distributed on an \"AS IS\" BASIS,\n * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.\n * See the License for the specific language governing permissions and\n * limitations under the License.\n */\n\n/* Modified by Nodus, 2026-11-09. Uncompiled native-ABI adaptation.\n * Derived from flashinfer::norm::RMSNormKernel, ordinary weight_bias=0 variant:\n * https://github.com/flashinfer-ai/flashinfer/blob/023a8474e1f7469572c0d4284535fc6ffa699574/include/flashinfer/norm.cuh\n * Original file SHA256: 31156798c949ae551b3149ce22c4edf80013b324a7ac695340dd378ea9903704\n * Changes: fixed BF16 ABI, 1-D block, scalar tail-safe column loops, explicit\n * CUDA intrinsics, static shared warp totals, and no host wrapper or PDL.\n * See derivation.json, LICENSE and NOTICE. This is not the upstream source.\n */\n\n/* Additional local fixture modification by Nodus, 2026-11-09.\n * Deliberately altered for a public-contract repair prompt preflight.\n * Not compiled, evaluated, or admitted as a training example.\n * See fixture-provenance.json for the complete authored change map.\n */\n\n#include <cuda_bf16.h>\n#include <cuda_runtime.h>\n\nextern \"C\" __global__ void nodus_rmsnorm_bf16(\n    const __nv_bfloat16* x, const __nv_bfloat16* w, __nv_bfloat16* out,\n    int T, int H, float eps) {\n  // The public ABI supplies grid=(T,1,1) and a whole-warp 1-D block.\n  // launch.json selects 256 threads; all live warps participate in reduction.\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 =\n      static_cast<unsigned long long>(blockIdx.x) * static_cast<unsigned long long>(H);\n 