CUDA SELU Activation Optimization Guide

__global__ void selu_basic(float* x, float* y, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float v = x[idx];
        y[idx] = 1.0507f * ((v > 0) ? v : 1.6733f * (expf(v) - 1.0f));
    }
}

__global__ void selu_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(
            SELU_LAMBDA * (v.x > 0 ? v.x : SELU_ALPHA * (__expf(v.x) - 1)),
            SELU_LAMBDA * (v.y > 0 ? v.y : SELU_ALPHA * (__expf(v.y) - 1)),
            SELU_LAMBDA * (v.z > 0 ? v.z : SELU_ALPHA * (__expf(v.z) - 1)),
            SELU_LAMBDA * (v.w > 0 ? v.w : SELU_ALPHA * (__expf(v.w) - 1))
        );
    }
}