CUDA Batch Exposure Processor

A GPU-accelerated image-processing project developed for the Johns Hopkins University CUDA at Scale for the Enterprise course on Coursera.

The program batch-processes TIFF aerial images using a custom CUDA kernel. Each GPU thread processes one image pixel and applies an exposure multiplier to its color channels.

OpenCV is used for TIFF image loading and saving, while the exposure operation itself is performed on the GPU using CUDA.

Features

• Batch processing of TIFF images
• Custom CUDA exposure kernel
• One CUDA thread per image pixel
• Configurable exposure multiplier
• CUDA error checking
• GPU kernel timing using CUDA events
• Automatic output directory generation
• Processing statistics
• Windows build and run scripts

Exposure Algorithm

Each color channel is multiplied by an exposure value:

output = input * exposure

The resulting value is clamped to the range:

0.0 to 1.0

For example:

Exposure 0.5 = darker
Exposure 1.0 = unchanged
Exposure 1.5 = brighter
Exposure 2.0 = brighter

The CUDA kernel performs this operation independently for every pixel in the image.

Processing Pipeline

TIFF Images
    |
    v
OpenCV Image Loading
    |
    v
Normalize Pixels to [0, 1]
    |
    v
Copy CPU -> GPU
    |
    v
CUDA Exposure Kernel
    |
    v
Copy GPU -> CPU
    |
    v
Convert to 8-bit Image
    |
    v
Save TIFF Output

Dataset

The program was tested using a collection of 38 aerial TIFF images of multiple resolutions.

A single execution processed:

Images processed: 38
Total pixels processed: 34,422,628

The input dataset contains approximately 101 MB of TIFF image data.

Rather than processing each image through a separate program execution, the application automatically finds and processes all .tif and .tiff files in the input directory.

Performance Results

The program was tested on:

GPU: NVIDIA GeForce RTX 4070
Exposure multiplier: 1.5

The recorded CUDA kernel results were:

Images processed: 38
Total pixels processed: 34,422,628
Total CUDA kernel time: 1.90272 ms
Average CUDA kernel time per image: 0.0500716 ms

CUDA events are used to measure the execution time of the exposure kernel.

These measurements represent CUDA kernel execution time only.

They do not include:

• TIFF image loading
• TIFF image saving
• GPU memory allocation
• CPU-to-GPU memory transfers
• GPU-to-CPU memory transfers
• Other CPU-side processing

Proof of Execution

The following screenshots show the aerial TIFF dataset before and after processing with the CUDA exposure kernel.

Original Images

The input dataset contains 38 aerial TIFF images.

CUDA Processed Images

The same dataset after applying an exposure multiplier of 1.5 using the custom CUDA kernel:

Terminal Results

The program successfully processed all 38 images using an NVIDIA GeForce RTX 4070.

The recorded execution processed:

Images processed: 38
Total pixels processed: 34,422,628

with the following CUDA kernel timing:

Total CUDA kernel time: 1.90272 ms
Average CUDA kernel time per image: 0.0500716 ms

Requirements

The project requires:

• NVIDIA CUDA-capable GPU
• NVIDIA CUDA Toolkit
• C++17
• OpenCV with TIFF support
• Microsoft Visual C++ Build Tools on Windows

The project was developed and tested with:

Operating System: Windows 11
GPU: NVIDIA GeForce RTX 4070
CUDA Toolkit: 13.2
OpenCV: 4.12.0
Compiler: Microsoft Visual C++ / nvcc

OpenCV was installed using vcpkg.

Project Structure

CUDA-Batch-Exposure-Processor/
|
|-- aerials/
|   |-- input TIFF images
|
|-- include/
|   |-- exposure.h
|
|-- src/
|   |-- exposure.cu
|   |-- main.cu
|
|-- output/
|   |-- generated exposure-adjusted TIFF images
|
|-- screenshots/
|   |-- before-processing.png
|   |-- after-processing.png
|   |-- terminal-results.png
|
|-- build.bat
|-- run.bat
|-- README.md
|-- .gitignore

Building the Project

Easy Build

A Windows build script is included with the project.

From the project root, run:

build.bat

The script compiles the CUDA source files and creates:

cuda_exposure.exe

Manual Build

On Windows, initialize the Visual Studio 2022 C++ build environment if necessary:

call "C:\Program Files (x86)\Microsoft Visual Studio\2022\BuildTools\VC\Auxiliary\Build\vcvars64.bat"

Verify that the Microsoft C++ compiler is available:

where cl

Verify that the CUDA compiler is available:

where nvcc

The project can then be compiled manually using:

nvcc -std=c++17 -gencode arch=compute_89,code=sm_89 src\main.cu src\exposure.cu ^
-Iinclude ^
-ID:\Documents\vcpkg\installed\x64-windows\include\opencv4 ^
-LD:\Documents\vcpkg\installed\x64-windows\lib ^
-lopencv_core4 ^
-lopencv_imgcodecs4 ^
-lopencv_imgproc4 ^
-o cuda_exposure.exe

The OpenCV paths shown above correspond to the development system and may need to be changed depending on the local vcpkg installation.

The compute_89 / sm_89 target is used for the NVIDIA GeForce RTX 4070. Users with different NVIDIA GPU architectures may need to change the CUDA architecture target.

Running the Program

Easy Run

A Windows run script is included.

To run using the default exposure multiplier of 1.5:

run.bat

To specify a different exposure multiplier:

run.bat 2.0

For example:

run.bat 0.5
run.bat 1.0
run.bat 1.5
run.bat 2.0

Run the Executable Directly

The program can also be executed directly:

cuda_exposure.exe 1.5

The command-line argument controls the exposure multiplier.

An exposure value greater than 1.0 brightens the image, while a value below 1.0 darkens it.

Input Images

Place TIFF images inside:

aerials/

The program automatically searches this directory for:

*.tif
*.tiff

Each valid TIFF image is processed during the same program execution.

Output Images

Processed images are automatically written to:

output/

For example:

aerials/2.1.01.tiff

produces:

output/2.1.01_exposure.tiff

The output directory is created automatically if it does not already exist.

CUDA Kernel Design

The custom exposure kernel uses a two-dimensional CUDA grid.

Each CUDA thread corresponds to one pixel in the image.

The pixel coordinates are calculated using:

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;

Threads outside the image dimensions are prevented from accessing invalid memory using a bounds check:

if (x >= width || y >= height)
{
    return;
}

The location of the pixel in the image buffer is calculated from its two-dimensional coordinates:

int index = (y * width + x) * channels;

The exposure operation is then applied to the pixel's color channels.

Conceptually:

output[channel] = clamp(input[channel] * exposure)

The resulting values are clamped to [0.0, 1.0] before the processed image is copied back to CPU memory.

The kernel uses a 16 x 16 CUDA thread block:

dim3 threadsPerBlock(16, 16);

The number of blocks is calculated from the dimensions of each input image so that the CUDA grid covers the complete image.

CPU and GPU Responsibilities

OpenCV handles CPU-side image operations:

TIFF decoding
Image format conversion
TIFF encoding

CUDA handles the exposure calculation:

Normalized Image
       |
       v
CPU -> GPU
       |
       v
ExposureKernel
       |
       v
GPU -> CPU
       |
       v
Processed Image

This separates image file handling from the GPU image-processing operation.

CUDA Memory Management

For each image, GPU memory is allocated for both the input and output image using:

cudaMalloc(...)

The normalized image is transferred from CPU memory to GPU memory using:

cudaMemcpy(
    ...,
    cudaMemcpyHostToDevice);

The CUDA kernel then processes the image.

After the kernel finishes, the processed image is transferred back to CPU memory using:

cudaMemcpy(
    ...,
    cudaMemcpyDeviceToHost);

The GPU allocations are then released using:

cudaFree(...)

CUDA Error Checking

CUDA operations are checked for errors throughout the program.

The program checks operations including:

cudaMalloc
cudaMemcpy
cudaMemset
CUDA kernel launch
CUDA event operations

Kernel launch errors are checked using:

cudaGetLastError()

CUDA error checking was useful during development for identifying configuration and execution problems.

GPU Timing

CUDA events are used to measure the execution time of the exposure kernel.

Timing begins immediately before the kernel launch and stops after the kernel completes.

The program reports the kernel execution time for each image and calculates the total and average kernel time across the entire dataset.

Example final output:

Processing complete.
GPU: NVIDIA GeForce RTX 4070
Images processed: 38
Total pixels processed: 34422628
Total CUDA kernel time: 1.90272 ms
Average CUDA kernel time per image: 0.0500716 ms

These values measure the CUDA kernel itself rather than the total application execution time.

Challenges

One challenge was handling TIFF image formats correctly.

The program loads TIFF images using cv::IMREAD_UNCHANGED so that the original image information is preserved. Supported images are then converted into normalized floating-point pixel data before being transferred to the GPU.

Another challenge was debugging the CPU-to-GPU image-processing pipeline. Pixel-value ranges were printed at different stages to verify that image data remained valid during processing:

Original TIFF
    |
    v
Normalized CPU Image
    |
    v
CUDA Kernel
    |
    v
GPU Output
    |
    v
Saved TIFF

CUDA toolkit and driver compatibility was another issue encountered during development.

The CUDA 13.2 compiler initially produced PTX that could not be executed by the installed NVIDIA driver. The project was instead compiled for the native architecture of the RTX 4070 using:

-gencode arch=compute_89,code=sm_89

This allowed the CUDA exposure kernel to execute successfully.

Integrating OpenCV with CUDA on Windows also required configuring the correct OpenCV include, library, and runtime DLL paths.

What I Learned

This project reinforced several CUDA and parallel-programming concepts:

• Mapping GPU threads to individual image pixels
• Working with two-dimensional CUDA grids
• Calculating CUDA grid and block dimensions
• Writing and launching custom CUDA kernels
• Performing bounds checking inside GPU kernels
• Allocating and freeing GPU memory
• Copying data between host and device memory
• Debugging CUDA runtime errors
• Measuring GPU kernel performance with CUDA events
• Integrating CUDA with OpenCV
• Processing multiple images during a single execution

The project also demonstrated the distinction between GPU computation and supporting CPU operations. OpenCV performs image file I/O and format conversion, while the custom CUDA kernel performs the exposure adjustment.

Project Purpose

The purpose of this project is to demonstrate image processing at scale using CUDA.

Instead of applying a GPU operation to a single test image, the application batch-processes an entire directory of TIFF images during one execution.

In the recorded run, the program processed:

38 TIFF images
34,422,628 total pixels

using a custom CUDA exposure kernel on an NVIDIA GeForce RTX 4070.

The project demonstrates a complete GPU image-processing workflow including image loading, host-to-device memory transfer, parallel CUDA execution, device-to-host transfer, output generation, CUDA error checking, and GPU kernel performance measurement.