Sigmoid Linear Unit

1
#include <cuda_runtime.h>
2

3
__global__ void silu_kernel(const float* input, float* output, int N) {
4
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
5
    if (idx > N) return;
6
    float sigma = 1.0 / (1.0 + exp(-input[idx]));
7
    float ans = input[idx] * sigma;
8
    output[idx] = ans;
9
}
10

11
// input, output are device pointers
12
extern "C" void solve(const float* input, float* output, int N) {
13
    int threadsPerBlock = 256;
14
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
15

16
    silu_kernel<<<blocksPerGrid, threadsPerBlock>>>(input, output, N);
17
    cudaDeviceSynchronize();
18
}