Swish-Gated Linear Unit

1
#include <cuda_runtime.h>
2

3
__global__ void swiglu_kernel(const float* input, float* output, int halfN) {
4
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
5
    int half = N / 2;
6
    if (idx >= N) return;
7
    float x1 = input[idx];
8
    float x2 = input[idx + half];
9
    float sigma = 1.0 / (1.0 + exp(-x1))
10
    float silu = x1 * sigma;
11
    float swiglu = silu * x2;
12
}
13

14
// input, output are device pointers
15
extern "C" void solve(const float* input, float* output, int N) {
16
    int halfN = N / 2;
17
    int threadsPerBlock = 256;
18
    int blocksPerGrid = (halfN + threadsPerBlock - 1) / threadsPerBlock;
19

20
    swiglu_kernel<<<blocksPerGrid, threadsPerBlock>>>(input, output, halfN);
21
    cudaDeviceSynchronize();
22
}