/**
 * Bias-add kernels for the batched cuBLAS forward pass.
 *
 * All-f32 pipeline: both output and bias are f32 (master weights).
 * No bf16 clamping needed — f32 has sufficient dynamic range.
 *
 * Launch config: grid=(ceil(total_elements/256), 1, 1), block=(256, 1, 1).
 */

/* f32 bias + ReLU for hidden layers (GemmEx TF32 → f32 output + f32 bias). */
extern "C" __global__ void add_bias_relu_f32_kernel(
    float* __restrict__ output,
    const float* __restrict__ bias,
    int out_dim,
    int total_elements)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= total_elements) return;
    float val = output[i] + bias[i % out_dim];
    output[i] = (val > 0.0f) ? val : 0.0f;  /* ReLU */
}

/* f32 bias (no activation) for logit output layers. */
extern "C" __global__ void add_bias_f32_kernel(
    float* __restrict__ output,
    const float* __restrict__ bias,
    int out_dim,
    int total_elements)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= total_elements) return;
    output[i] = output[i] + bias[i % out_dim];
}
