{
  "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  __shared__ float warp_sums[32];\n\n  // Scalar specialization of the original vector accumulation, including tails.\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  // Preserve the source kernel's warp XOR reduction, then first-warp reduction.\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"
}<|im_end|>