1
#include <cuda_runtime.h>
2

3
__global__ void invert_kernel(unsigned char* image, int width, int height) {
4
    int x = blockDim.x * blockIdx.x + threadIdx.x;
5
    int id = x * 4;
6
    if (id < width * height * 4) {
7
        image[id] = 255 - image[id];
8
        image[id + 1] = 255 - image[id + 1];
9
        image[id + 2] = 255 - image[id + 2];
10
        image[id + 3] = image[id + 3];
11
    }
12
}
13
// image_input, image_output are device pointers (i.e. pointers to memory on the GPU)
14
extern "C" void solve(unsigned char* image, int width, int height) {
15
    int threadsPerBlock = 256;
16
    int blocksPerGrid = (width * height + threadsPerBlock - 1) / threadsPerBlock;
17

18
    invert_kernel<<<blocksPerGrid, threadsPerBlock>>>(image, width, height);
19
    cudaDeviceSynchronize();
20
}