/* SPDX-License-Identifier: AGPL-3.0-or-later */
/* Copyright (C) 2026 David (Corey) Herrera */
/*
 * elementwise.cu — CUDA elementwise kernels for Foundry
 *
 * Target: sm_37 (K80, Kepler)
 * Compile: nvcc -ptx -arch=sm_37 elementwise.cu -o elementwise.ptx
 *
 * Constraints:
 *   - No FP16 math (sm_37 has no native half support)
 *   - Use __shfl_down() not __shfl_down_sync() (sync is sm_70+)
 *   - Max 1024 threads per block, 48 KB shared memory per SM
 */

extern "C" {

/* out[i] = a[i] + b[i] */
__global__ void fc_add(const float *a, const float *b, float *out, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        out[idx] = a[idx] + b[idx];
    }
}

/* out[i] = a[i] * b[i] */
__global__ void fc_mul(const float *a, const float *b, float *out, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        out[idx] = a[idx] * b[idx];
    }
}

/* out[i] = x[i] * sigmoid(x[i]) = x[i] / (1 + exp(-x[i])) */
__global__ void fc_silu(const float *x, float *out, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float v = x[idx];
        out[idx] = v / (1.0f + expf(-v));
    }
}

} /* extern "C" */