Leaky ReLU

1
#include <cuda_runtime.h>
2

3
__global__ void leaky_relu_kernel(const float* input, float* output, int N) {
4
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
5
    if (idx < N) {
6
        output[idx] = input[idx] > 0 ? input[idx] : 0.01 * input[idx];
7
    }
8
}
9

10
// input, output are device pointers (i.e. pointers to memory on the GPU)
11
extern "C" void solve(const float* input, float* output, int N) {
12
    int threadsPerBlock = 256;
13
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
14

15
    leaky_relu_kernel<<<blocksPerGrid, threadsPerBlock>>>(input, output, N);
16
    cudaDeviceSynchronize();
17
}