Assignment 4 : Image Filters Using Cuda: Exact Answer & Steps

Assignment 4: Image Filters Using CUDA

So you're tackling Assignment 4 — building image filters with CUDA. If you're like most students approaching this, you probably have two reactions: excitement that you're finally getting to do some real GPU programming, and maybe a little panic because parallelizing image processing has some tricky parts that aren't obvious until you actually write the code.

That's exactly why I'm writing this. Not to do your assignment for you, but to walk through the concepts, show you what actually works, and help you avoid the mistakes that make this assignment take three times longer than it should.

Let's dig in.

What Is CUDA Image Filtering

At its core, image filtering with CUDA means using the GPU to apply mathematical operations to every pixel in an image — simultaneously, not one at a time. A filter is just a small matrix (usually 3x3 or 5x5) that you slide across the entire image, computing a new value for each pixel based on its neighbors.

Think about what that means sequentially: for a 1920x1080 image, you're doing roughly 2 million pixels times however many operations your filter needs. Practically speaking, on a CPU, that's a loop that runs millions of times. On a GPU with CUDA, you can launch thousands of threads where each thread handles one pixel.

The filters you'll probably implement include box blur (simple average of neighbors), Gaussian blur (weighted average that looks more natural), Sobel edge detection (finds edges by comparing gradients), and maybe sharpen (enhances contrast between a pixel and its neighbors).

The key insight is this: each pixel's new value is independent of other pixels' new values. That's what makes it embarrassingly parallel — perfect for GPU execution.

The CUDA Thread Model for Images

Here's how it maps. That said, your image is a 2D grid. So in CUDA, you'll organize your threads into a 2D grid of blocks. Each thread calculates the output for one pixel.

The usual setup looks something like this:

dim3 blockSize(16, 16);
dim3 gridSize((width + blockSize.x - 1) / blockSize.x,
              (height + blockSize.y - 1) / blockSize.y);
kernelName<<>>(...);

That gives you 256 threads per block, which is a solid starting point. Each thread knows its pixel coordinates from blockIdx and threadIdx, and it calculates where in global memory that pixel lives.

Memory Access Patterns Matter

This is where things get real. Think about it: global memory on the GPU is slow — much slower than the compute units. In real terms, when threads in a warp (32 threads that execute together) access global memory, you want them to access consecutive addresses. That's called coalesced memory access.

For image processing, if your threads read from row-major order in memory (which is standard), you're probably fine for the main pixel. But if you're reading from a filter kernel — the small matrix of weights — you might have threads in the same warp reading from scattered addresses. That's when you move data to shared memory.

Why It Matters

Why does this assignment matter beyond the grade? Because you're learning skills that apply everywhere in high-performance computing.

The GPU is a massively parallel processor. Image filtering is one of the clearest examples of parallelizable work.Master this, and you understand the core ideas behind:

• Real-time graphics processing — every Instagram filter, every video call enhancement, everything running in Photoshop — all of it uses GPU parallelism
• Deep learning inference — those neural networks processing images? Convolution operations look a lot like what you're doing in this assignment
• Scientific computing — PDE solvers, fluid simulations, molecular dynamics all use the same patterns

Also, let's be honest: it's going to be visibly faster. A naive CUDA implementation will absolutely outperform a naive CPU implementation, and a well-optimized one will leave the CPU in the dust. Seeing your code run hundreds of times faster on the GPU than the equivalent CPU version is genuinely satisfying.

When CUDA Beats CPU

The rule of thumb: if your operation is parallelizable and you're processing enough data, GPU wins. For image filters on typical image sizes (megapixels, not pixels), CUDA wins.

Where it doesn't make sense: tiny images, or filters so simple the overhead of transferring data from CPU to GPU dominates. That's why you'd never use CUDA to apply a filter to a 64x64 icon.

How It Works

Here's the step-by-step of building a CUDA image filter, assuming you're implementing a convolution-based filter like blur or edge detection.

Step 1: Allocate Memory

You need memory on both the host (CPU) and device (GPU).

// Host memory
unsigned char *h_input, *h_output;
// Device memory
unsigned char *d_input, *d_output;

cudaMalloc(&d_input, width * height * channels);
cudaMalloc(&d_output, width * height * channels);

Don't forget to check for errors on your CUDA calls. It's good practice, and it'll save you hours of debugging when something goes wrong.

Step 2: Copy Data to Device

cudaMemcpy(d_input, h_input, width * height * channels, cudaMemcpyHostToDevice);

This is where pinned memory can help. If you allocate host memory as page-locked (using cudaMallocHost), transfers are faster. For Assignment 4, it's probably optional, but it's worth knowing.

Step 3: Launch the Kernel

Your kernel is the function that runs on the GPU. Each thread computes one output pixel.

__global__ void convolveKernel(unsigned char *input, unsigned char *output,
                               int width, int height, float *kernel, int kSize) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    
    if (x >= width || y >= height) return;
    
    // Apply convolution for this pixel
    float sum = 0.0f;
    int half = kSize / 2;
    
    for (int ky = -half; ky <= half; ky++) {
        for (int kx = -half; kx <= half; kx++) {
            int px = x + kx;
            int py = y + ky;
            
            // Boundary check - skip or clamp
            if (px >= 0 && px < width && py >= 0 && py < height) {
                int pixelIdx = (py * width + px);
                float weight = kernel[(ky + half) * kSize + (kx + half)];
                sum += input[pixelIdx] * weight;
            }
        }
    }
    
    output[y * width + x] = (unsigned char)sum;
}

This is the basic structure. You'll need to handle channels (RGB) separately or process them together depending on your filter.

Step 4: Copy Results Back

cudaMemcpy(h_output, d_output, width * height * channels, cudaMemcpyDeviceToHost);

Step 5: Clean Up

cudaFree(d_input);
cudaFree(d_output);
// Free host memory too

Optimization: Using Shared Memory

The code above works, but it reads from global memory for every neighbor pixel. Shared memory is faster. The idea: load the region each block needs into shared memory, then have all threads in the block read from that fast on-chip memory.

Continue exploring with our guides on world war 2 map activity and which tineco vacuum is the best.

__shared__ unsigned char sharedImage[BLOCK_SIZE][BLOCK_SIZE];

The tricky part is handling the edges — threads at the block boundary need to read neighbor pixels that are outside the block. The common approach is to have each thread load one pixel, then synchronize, then do the computation. Or load extra pixels at the edges.

This is where the real learning happens. Getting shared memory right is the difference between "works" and "fast."

Common Mistakes

Here's where students get stuck. I've seen these issues repeatedly.

Boundary Conditions

The most common bug: your kernel tries to read pixels outside the image. For pixels near the edge, you need to decide what to do. Options include:

• Clamping — pretend the edge pixel extends infinitely
• Wrapping — wrap around to the other side
• Ignoring — just don't process edge pixels (but then your output is smaller)

The code I showed above checks bounds inside the loop. That's correct but slow — every thread checks bounds for every neighbor. A cleaner approach: handle the main area with one kernel, edges with another, or use a border of dummy values.

Thread Indexing Errors

Getting x and y wrong is easy. Remember:

• blockIdx.x is which block horizontally
• threadIdx.x is which thread inside that block
• So blockIdx.x * blockDim.x + threadIdx.x gives the global position

Off-by-one errors here mean your output is shifted, cropped, or garbage.

Forgetting to Synchronize

If you're using shared memory, __syncthreads() is your friend. Every thread in a block must hit the synchronization point before any can continue. Miss it and you'll read data that hasn't been written yet — random garbage, usually.

Not Using the Right Kernel

CUDA has different kernel versions. Which means make sure you're calling the right one with the right parameters. Silly mistake, but it happens. The details matter here.

Ignoring Memory Transfer Overhead

If you're processing a tiny image, the time to copy data to the GPU and back can exceed the time to just do it on the CPU. For Assignment 4, you're probably using a decent-sized image, but it's worth knowing.

Practical Tips

Here's what actually works, beyond the textbook approach.

Start simple. Get a basic box blur working before you tackle Gaussian or edge detection. Get the memory, kernel launch, and copy-back working first. Then optimize.

Use cudaDeviceSynchronize() for debugging. Put it after your kernel launch and check cudaGetLastError(). It'll tell you if your kernel crashed and why — out of memory, invalid configuration, something else.

Print strategically. printf works in CUDA kernels (for compute capability 2.0+). It's incredibly useful for seeing what's actually happening.

Block size of 16x16 is a safe default. But you can experiment. Some filters benefit from 32x1 or other shapes depending on memory access patterns.

Think about memory coalescing. Threads in a warp (32 consecutive threads) should read consecutive addresses. For row-major images, threads reading the same row are coalesced. That's good.

For separable filters like Gaussian, use two passes. A 2D Gaussian blur can be done as two 1D blurs (horizontal then vertical). It's much faster. If your assignment allows, this is a huge win.

FAQ

How do I handle RGB channels?

For each pixel, you have three values (or four with alpha). That's why the simplest approach: process them identically, applying the same filter to R, G, and B separately. Some filters might treat channels differently, but for basic blur and sharpen, independent channel processing works fine.

What's the difference between global and shared memory?

Global memory is on the GPU board — gigabytes, but slow access (hundreds of cycles). Shared memory is on each streaming multiprocessor — kilabytes, but extremely fast (single-digit cycles). Shared memory is like a cache you control manually.

Why is my output wrong at the edges?

Probably missing boundary handling. The filter kernel extends beyond the image for edge pixels. Add checks or use a padded input image.

How do I know if my CUDA code is actually running in parallel?

Watch the GPU utilization with nvidia-smi while your code runs. If it's high (70%+), your code is using the GPU. If it's near zero, something's wrong — probably a synchronous issue or your kernel isn't being launched.

Should I use texture memory?

Texture memory provides hardware interpolation and caching that's useful for some filters. And for simple convolution, it's probably overkill. But if you're doing something like zoom or rotation, texture memory helps.

Closing

CUDA image filters are one of those assignments that teach you more than just the specific topic. You're learning parallel thinking, memory optimization, and how to debug code that runs on thousands of threads simultaneously.

Start with a working naive implementation. Get it correct first. Here's the thing — then profile, identify bottlenecks, and optimize. That's the real workflow — not writing perfect code from the start, but iterating toward better code.

You've got this.