CUDA Tanh Activation Optimization Guide

__device__ float tanh_manual(float x) {
    float ep = expf(x), em = expf(-x);
    return (ep - em) / (ep + em);  // Two exp calls!
}

// In LSTM: gates use sigmoid, cell state uses tanh
// Fuse all gate computations
__global__ void lstm_gates_fused(float4* gates, float4* output, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float4 g = gates[idx];
        // i, f, o gates: sigmoid; g gate: tanh
        output[idx] = make_float4(
            sigmoid_stable(g.x),  // input gate
            sigmoid_stable(g.y),  // forget gate
            sigmoid_stable(g.z),  // output gate
            tanhf(g.w)            // cell gate
        );
    }
}