CUDA Swish/SiLU Activation Optimization Guide

__global__ void swish_naive(float* x, float* y, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float v = x[idx];
        float sig = 1.0f / (1.0f + expf(-v));
        y[idx] = v * sig;
    }
}

__global__ void swish_opt(float4* x, float4* y, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float4 v = x[idx];
        y[idx] = make_float4(
            swish_stable(v.x), swish_stable(v.y),
            swish_stable(v.z), swish_stable(v.w)
        );
    }
}

__device__ float swish_stable(float x) {
    float ax = fabsf(x);
    float ex = __expf(-ax);
    return (x >= 0) ? x / (1.0f + ex) : x * ex / (1.0f + ex);
}