Value Clipping

1
#include <cuda_runtime.h>
2

3
__global__ void clip_kernel(const float* input, float* output, float lo, float hi, int N) {
4
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
5
    if (idx >= N) return;
6
    if (input[idx] >= lo && input[idx] <= hi) {
7
        output[idx] = input[idx];
8
    } else if (input[idx] < lo) {
9
        output[idx] = lo;
10
    } else {
11
        output[idx] = hi;
12
    }
13
}
14

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

20
    clip_kernel<<<blocksPerGrid, threadsPerBlock>>>(input, output, lo, hi, N);
21
    cudaDeviceSynchronize();
22
}