G2ES GPU Image Processing Pipeline

A CUDA-accelerated image processing pipeline that implements four classic computer vision stages — Grayscale, Gaussian blur, Equalization (histogram), Sobel edge detection — with a built-in CPU reference implementation for benchmarking.

📖 中文文档请查看 README_CN.md

🎨 Pipeline Demo

Input	Grayscale	Gaussian Blur	Histogram Eq.	Sobel Edge

📑 Table of Contents

• 🎨 Pipeline Demo
• 📋 Overview
• 📊 Performance Overview
• 🔄 Pipeline Stages
• 📁 Project Structure
• ⚙️ Prerequisites
• 🔨 Build
• 🚀 Usage
• 📤 Output Files
• 🏗️ Architecture Details
• 🧯 Error Handling
• 📊 Performance Benchmarking
• 📄 License

📋 Overview

G2ES loads an image (any format OpenCV supports: PNG, JPG, BMP, PGM, PPM, etc.), processes it through a four-stage pipeline, and writes intermediate and final results as PNG files. The project supports four execution modes:

Mode	Flag	Description
GPU only	--gpu (default)	Runs the CUDA pipeline with naive kernels
GPU optimized	--gpu-optimized	Runs the CUDA pipeline with optimized kernels
CPU only	--cpu	Runs the single-threaded CPU reference pipeline
Both	--both	Runs both CPU and GPU (naive) and prints a speedup comparison

📊 Performance Overview

Measured on Jetson AGX Orin with a 1440×810 image.

Metric	CPU	GPU (Naive)	GPU Optimized (No Sobel Opt) 🏆	GPU Optimized (All)
Kernel compute	34.29 ms	3.31 ms	1.66 ms	1.89 ms
+ Transfer (H2D+D2H)	—	6.30 ms	4.60 ms	4.85 ms
Speedup vs CPU (kernel only)	1×	10.4×	20.7×	18.1×
Speedup vs CPU (incl. transfer)	—	5.4×	7.5×	7.1×

Key takeaways:

• 🏆 Best config = optimized RGB + optimized Gaussian + optimized Histogram + naive Sobel → 20.7× kernel speedup
• Full optimization (1.89 ms) is 13.8% slower than the "no Sobel opt" variant (1.66 ms)—the shared memory Sobel backfires (see Shared Memory Optimization Insights)
• Data transfers (~3 ms) remain the end-to-end bottleneck, capping the transfer-inclusive speedup at 7.5×

🔄 Pipeline Stages

Input Image → [1. RGB→Gray] → [2. Gaussian Blur] → [3. Histogram Eq.] → [4. Sobel Edge] → Output

Stage	Algorithm	Kernel Size	Key Detail
1. RGB to Grayscale	Weighted luminance 0.299R + 0.587G + 0.114B	—	ITU-R BT.601 standard
2. Gaussian Blur	5×5 Gaussian smoothing	5×5	Weights stored in __constant__ memory
3. Histogram Equalization	CDF-based intensity redistribution	—	3-kernel decomposition (histogram → CDF/LUT → apply)
4. Sobel Edge Detection	min(255, √(Gx² + Gy²))	3×3	Naive implementation (baseline)

Optimized Kernels

The --gpu-optimized mode uses optimized kernel implementations that provide better performance:

Stage	Optimization Technique	Measured Performance Impact
1. RGB to Grayscale	uchar3 vectorized memory access	+3.2%
2. Gaussian Blur	Shared memory + separable convolution	+1.7%
3. Histogram	Shared memory local histograms + grid-stride loop	+87.8% 🏆
4. Sobel Edge	Shared memory tile with halo	-96.6% ❌

Note: The above figures are measured on a 1440×810 image on Jetson AGX Orin. Performance varies by hardware, image size, and implementation details. The Sobel regression is explained in the Shared Memory Optimization Insights section below.

📁 Project Structure

G2ES_GPU_Pipeline/
├── include/
│   ├── kernels.cuh            # All __global__ kernel declarations
│   ├── pipeline_common.h      # Pipeline common definitions and structures
│   └── utils.h                # CUDA error-checking macro
├── src/
│   ├── main.cu                # Entry point, CLI parsing, benchmark orchestration
│   ├── cpu_pipeline.cpp       # CPU pipeline implementation
│   ├── cpu_pipeline.h         # CPU pipeline header
│   ├── gpu_pipeline.cu        # GPU pipeline implementation (naive kernels)
│   ├── gpu_pipeline.h         # GPU pipeline header
│   ├── gpu_pipeline_optimized.cu  # GPU pipeline implementation (optimized kernels)
│   ├── pipeline_common.cu     # Pipeline common utilities
│   └── kernels/
│       ├── rgb_to_gray.cu     # RGB → Grayscale kernel
│       ├── gaussian_blur.cu   # 5×5 Gaussian blur kernel
│       ├── histogram_equalization.cu  # Histogram + CDF/LUT + LUT apply kernels
│       └── sobel_edge.cu      # Sobel edge detection kernel
├── image/
│   └── test_image.png         # Sample input image
├── Makefile                   # nvcc-based build system
├── README.md                  # This file
└── README_CN.md               # 中文文档

⚙️ Prerequisites

Dependency	Version	Purpose
NVIDIA CUDA Toolkit	11.0+	nvcc compiler and CUDA runtime
OpenCV	4.x	Image I/O (imread/imwrite) and color conversion
CUDA-capable GPU	Compute Capability 8.7	Default target: Jetson AGX Orin / Ampere

Note: The Makefile uses -arch=sm_87 by default. If your GPU has a different compute capability, modify the NVCCFLAGS in the Makefile accordingly (e.g., sm_80 for A100, sm_86 for RTX 3080, sm_89 for RTX 4090).

Installing Dependencies (Ubuntu/Debian)

# CUDA Toolkit (if not already installed)
sudo apt install nvidia-cuda-toolkit

# OpenCV 4
sudo apt install libopencv-dev

# Verify
nvcc --version
pkg-config --modversion opencv4

🔨 Build

make

This compiles all .cu source files under src/ and src/kernels/ into object files in obj/, then links them into the image_pipeline binary.

To clean build artifacts:

make clean

🚀 Usage

./image_pipeline [mode] <input_image> <output_prefix>

Examples

# GPU mode (default) — reads test_image.png, writes test_image_*.png
./image_pipeline image/test_image.png image/test_image

# GPU optimized mode — uses optimized kernels for better performance
./image_pipeline --gpu-optimized image/test_image.png image/test_image_optimized

# CPU mode only
./image_pipeline --cpu image/test_image.png image/test_image_cpu

# Run both GPU and CPU, print speedup comparison
./image_pipeline --both image/test_image.png image/test_image

Arguments

Argument	Required	Description
--gpu / --gpu-optimized / --cpu / --both	No	Execution mode (default: --gpu)
<input_image>	Yes	Path to input image (any OpenCV-supported format)
<output_prefix>	Yes	Prefix for output file names

📤 Output Files

Each run produces four PNG images per mode:

File	Description
{prefix}_gray.png	Grayscale conversion result
{prefix}_blur.png	Gaussian blur result
{prefix}_equalized.png	Histogram equalization result
{prefix}_edge.png	Sobel edge detection result

In --both mode, files are automatically suffixed with _cpu or _gpu (e.g., test_image_cpu_gray.png, test_image_gpu_gray.png).

🏗️ Architecture Details

GPU Kernel Design

Thread Block Layout:

• Image-processing kernels (grayscale, blur, sobel): 2D blocks of 16×16 = 256 threads
• Linear kernels (histogram, LUT apply): 1D blocks of 256 threads

Memory Strategy:

• Intermediate buffers (d_gray, d_blur, d_equalized, d_edge) stay on the GPU throughout the pipeline — only the final result is copied back to the host
• Gaussian kernel weights use __constant__ memory for broadcast efficiency
• Histogram equalization uses shared memory for parallel prefix sum (Blelloch scan)

Histogram Equalization Decomposition:

The histogram equalization stage is split into three specialized kernels:

1. compute_histogram_global_kernel — Each thread processes one pixel, uses atomicAdd to build a 256-bin histogram in global memory
2. build_equalization_lut_kernel — Single block of 256 threads performs a Blelloch inclusive scan in shared memory to compute the CDF, then builds the equalization LUT
3. apply_lut_kernel — Each thread does a simple table lookup: output[i] = lut[input[i]]

CPU Reference Implementation

The CPU pipeline in src/cpu_pipeline.cpp implements identical algorithms using single-threaded nested loops. It exists for:

• Correctness verification (comparing GPU output against CPU output)
• Performance benchmarking (measuring GPU speedup)

🧯 Error Handling

All CUDA API calls are checked through the CHECK_CUDA_ERROR macro defined in include/utils.h. On failure, the helper prints detailed error information, including file name, line number, CUDA call, and error code, then terminates the program.

📊 Performance Benchmarking

Use --both mode to compare GPU and CPU performance in a single run:

./image_pipeline --both image/test_image.png image/test_image

The output includes:

• Per-stage timing for both CPU and GPU
• Total compute time
• Kernel-only speedup (GPU kernel execution vs CPU compute)
• Transfer-inclusive speedup (including host↔device memory transfers)

GPU timing uses cudaEvent for accurate kernel measurement; CPU timing uses std::chrono::steady_clock.

Comparing Naive vs Optimized Kernels

To compare the performance of naive and optimized GPU kernels:

# Run naive kernels
./image_pipeline --gpu image/test_image.png output/naive

# Run optimized kernels
./image_pipeline --gpu-optimized image/test_image.png output/optimized

Example Performance Results (1440×810 image on Jetson AGX Orin):

Metric	Naive Kernels	Optimized Kernels	Improvement
Kernel Total	3.31 ms	1.89 ms	42.9% faster
Transfer+Kernel	6.30 ms	4.85 ms	23.1% faster

💡 The full CPU vs GPU vs GPU Optimized comparison table (including the best-performing "No Sobel Opt" variant at 20.7× speedup) is in the Performance Overview at the top of this document.

Kernel-by-Kernel Benchmark

Use --benchmark mode to test each optimized kernel individually with controlled variables:

./image_pipeline --benchmark image/test_image.png

This runs 6 tests (3 warmup + 10 benchmark runs each):

1. Baseline (all naive kernels)
2. Only RGB optimized
3. Only Gaussian Blur optimized
4. Only Histogram optimized
5. Only Sobel optimized
6. All optimized

Example Benchmark Results (1440×810 image on Jetson AGX Orin):

Test Configuration	Avg Time	Min Time	Improvement
Baseline (All Naive)	0.6518 ms	0.6463 ms	-
Only RGB Optimized	0.6544 ms	0.6444 ms	-0.40%
Only Gaussian Blur Optimized	0.6699 ms	0.6623 ms	-2.78%
Only Histogram Optimized	0.2882 ms	0.2770 ms	+55.78%
Only Sobel Optimized	0.7049 ms	0.6973 ms	-8.15%
All Optimized	0.3366 ms	0.3306 ms	+48.36%

Key Findings:

• Histogram optimization provides the most significant improvement (55.78%)
• Other individual optimizations show slight overhead
• Combined optimization achieves 48.36% improvement, mainly from histogram optimization

Shared Memory Optimization Insights

Shared memory is one of the most commonly used CUDA optimization techniques. The basic idea is to cooperatively load data from global memory into shared memory (on-chip SRAM), then have threads exchange data through shared memory, reducing redundant requests for global memory bandwidth.

Using the Sobel 3×3 edge detection as a case study, the optimized version (sobel_edge_optimized_kernel) employs a shared memory tile with halo:

1. Each 16×16 thread block cooperatively loads an 18×18 tile (16×16 output region + 1-pixel halo on each side) into shared memory
2. __syncthreads() barrier ensures all data is ready
3. Each thread reads the 3×3 neighborhood from shared memory and computes the Sobel gradient

Why is the shared memory Sobel 96.6% slower?

Factor	Explanation
🔄 Cooperative load overhead	Each block must load the halo region (18×18=324 elements vs 16×16=256 output), a 26.6% increase in load traffic
⏱ Sync barrier	__syncthreads() introduces warp-level waiting, which is significant at high occupancy
🎯 Cache hit rate	A 3×3 neighborhood has small stride; adjacent threads' reads already exhibit good spatial locality in L1/L2 cache. The naive version hits L1 cache in practice, making latency already quite low
🧮 Kernel size	The 3×3 kernel has extremely low arithmetic intensity (~2 FLOP/byte), leaving very little headroom for optimization gains

When does shared memory optimization actually help?

The effectiveness of shared memory depends on the trade-off between bandwidth savings and cooperative load + sync overhead:

Dimension	Shared memory helps ✅	Neutral ⚖️	Shared memory hurts ❌
Kernel size	Large (7×7+)	Medium (5×5)	Small (3×3)
Access span	Large stride	Moderate	Compact neighborhood
Cache quality	Small / slow L1	Average	Large / fast L1
Example in project	—	Gaussian Blur 5×5	Sobel 3×3

• Histogram (+87.8% ✅): Heavy atomicAdd contention on global memory — shared memory local reduction drastically reduces global atomic conflicts, yielding huge gains
• Gaussian Blur 5×5 (+1.7% ≈ neutral): The 5×5 kernel is at the break-even point — the cache already covers most redundant reads, leaving little for shared memory to improve
• Sobel 3×3 (-96.6% ❌): The kernel is too small — cooperative load + sync overhead exceeds whatever bandwidth savings the cache hasn't already delivered

Rule of thumb: For kernels with radius ≤ 2 (e.g., 3×3, 5×5), try the naive direct-read version first. If the L1 cache hit rate is already high, the cooperative load and synchronization overhead of shared memory may hurt performance. For radius ≥ 3 or strided access patterns, shared memory often brings clear benefits. Always rely on measurements — performance expectations need experimental validation, not intuition.

For detailed experimental records and discussion, see the inline comments in sobel_edge.cu (lines 73-88).

📄 License

This project is for educational/study purposes.