CUDA & GPU Vectorization — Technical Reference

Compiled from a full technical deep-dive session covering vectorization in databases, CUDA image processing APIs, CPU–GPU communication, and event monitoring internals.

CUDA APIs for Vectorized Query Processing
CPU SIMD vs Vulkan vs CUDA — Comparison
Vectorization in Open-Source Databases — Examples
Vectorization API Usage Per Database
CUDA Image Processing APIs
Legacy vs Modern CUDA Image APIs
CPU–GPU Communication: Events and Messages
CUDA Event Monitoring Internals: Polling vs Interrupts

1. CUDA APIs for Vectorized Query Processing

1.1 Vectorized Memory Access — Vector Types

Using vectorized loads reduces the total number of instructions, reduces latency, and improves bandwidth utilization. Use the vector data types defined in the CUDA C++ standard headers.

// Load 4 floats in a single instruction (128-bit)
float4 val = *reinterpret_cast<const float4*>(ptr + i);

// Load 4 ints in a single instruction
int4  ival = *reinterpret_cast<const int4*>(col + i);

Applicable query operators: column scans, hash table probing, equality filters.

1.2 CUB (CUDA UnBound)

CUB provides fast primitives for scan, sort, reduction, and device-side allocators. It is the building block library for query operators.

Primitive	Query use case
`cub::DeviceReduce`	Aggregation: SUM, MIN, MAX
`cub::DeviceScan`	Prefix sums for GROUP BY / compaction
`cub::DeviceSelect`	WHERE clause stream compaction
`cub::DeviceRadixSort`	Sort-merge joins, ORDER BY
`cub::WarpReduce`	Per-warp aggregation in 5 steps
`cub::BlockReduce`	Per-block aggregation

#include <cub/cub.cuh>

// Allocate temp storage
void   *d_temp = nullptr;
size_t  temp_bytes = 0;
cub::DeviceReduce::Sum(d_temp, temp_bytes, d_in, d_out, N);
cudaMalloc(&d_temp, temp_bytes);

// Run reduction
cub::DeviceReduce::Sum(d_temp, temp_bytes, d_in, d_out, N);

1.3 Thrust

Thrust provides high-level host/device vector containers and algorithms modelled on the C++ STL.

#include <thrust/device_vector.h>
#include <thrust/reduce.h>
#include <thrust/copy.h>

thrust::device_vector<float> d_col(h_col.begin(), h_col.end());

// Vectorized filter + reduction
float total = thrust::reduce(
    thrust::make_transform_iterator(d_col.begin(), pred_fn),
    thrust::make_transform_iterator(d_col.end(),   pred_fn));

1.4 Warp Shuffle — Warp-Level Reduction

// Sum a float across all 32 lanes of a warp
__device__ float warp_reduce_sum(float val) {
    for (int offset = 16; offset > 0; offset >>= 1)
        val += __shfl_down_sync(0xFFFFFFFF, val, offset);
    return val;  // lane 0 holds the total
}

__shfl_down_sync butterfly reduction completes in 5 steps for 32 lanes, entirely in registers with no shared memory traffic.

1.5 Warp Ballot — Predicate Vectorization

// Build a 32-bit bitmask of lanes satisfying the predicate
__global__ void filter_kernel(const int* col, int threshold,
                               int* out_count) {
    int val  = col[blockIdx.x * blockDim.x + threadIdx.x];
    int mask = __ballot_sync(0xFFFFFFFF, val > threshold);

    if (threadIdx.x == 0)
        atomicAdd(out_count, __popc(mask));  // count matching rows
}

__ballot_sync is the GPU analog of SIMD predicate vectorization on CPUs.

1.6 SWAR (SIMD Within A Register) + 128-bit Loads

128-bit loads return multiple 64-bit words in one instruction. Each word is compared against a broadcast value using constant-time arithmetic, eliminating branching.

// 128-bit load using int4 — reads 4 × 32-bit integers at once
__global__ void probe_hash_table(const int* keys, int target,
                                  int* match_count, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx * 4 >= N) return;

    int4 chunk = reinterpret_cast<const int4*>(keys)[idx];

    int hits = (chunk.x == target) + (chunk.y == target)
             + (chunk.z == target) + (chunk.w == target);
    atomicAdd(match_count, hits);
}

1.7 CUDA Streams + Async Memory Copy

Overlapping data transfer and kernel execution hides PCIe latency.

cudaStream_t s0, s1;
cudaStreamCreate(&s0);
cudaStreamCreate(&s1);

// Stream 0: transfer batch A while stream 1 processes batch B
cudaMemcpyAsync(d_A, h_A, bytes, cudaMemcpyHostToDevice, s0);
query_kernel<<<grid, block, 0, s1>>>(d_B, d_out_B, N);

cudaStreamSynchronize(s0);
cudaStreamSynchronize(s1);

1.8 cuFFT — Frequency-Domain Filtering

#include <cufft.h>

cufftHandle plan;
cufftPlan2d(&plan, height, width, CUFFT_R2C);

cufftExecR2C(plan, d_input, d_freq);
// apply filter in frequency domain ...
cufftExecC2R(plan, d_freq, d_output);

cufftDestroy(plan);

Used for: low-pass / high-pass filters, motion blur deconvolution, phase correlation image alignment.

1.9 API Summary Table

API / Feature	Query Use Case
`int4` / `float4` vector types	Coalesced column scans
`cub::DeviceSelect`	WHERE clause / filter compaction
`cub::DeviceScan`	Prefix sum for GROUP BY / aggregation
`cub::DeviceRadixSort`	Sort-merge joins, ORDER BY
Thrust algorithms	High-level query pipeline prototyping
`__shfl_down_sync`	Warp-level reduction for SUM/MIN/MAX
`__ballot_sync`	Bitmap predicate evaluation
SWAR + 128-bit loads	Hash table probing, equality filters
CUDA Streams	Overlapping scan batches with data transfer
cuBLAS/cuBLASLt	Vector similarity and analytical ML queries

2. CPU SIMD vs Vulkan vs CUDA — Comparison

2.1 Execution Model

	CPU SIMD / AVX	Vulkan Compute	CUDA
Model	SIMD — one instruction, packed register	SIMT — subgroup lanes in lockstep	SIMT — warp of 32 threads
Lane width	128-bit (SSE) / 256-bit (AVX2) / 512-bit (AVX-512)	8–128 lanes (vendor-defined)	Fixed 32 threads per warp
Portability	x86-only	Cross-vendor GPU (NVIDIA/AMD/Intel/ARM)	NVIDIA-only

2.2 Key Vectorization APIs

CPU SIMD (x86 intrinsics):

// AVX2 — add 8 floats in one instruction
__m256 a   = _mm256_load_ps(src_a);
__m256 b   = _mm256_load_ps(src_b);
__m256 res = _mm256_add_ps(a, b);

// AVX-512 — fused multiply-add on 16 floats
__m512 r = _mm512_fmadd_ps(a, b, c);

// Predicate → 8-bit bitmask
__m256 cmp  = _mm256_cmp_ps(a, threshold, _CMP_GT_OQ);
int    mask = _mm256_movemask_ps(cmp);

Vulkan (GLSL compute shader):

#extension GL_KHR_shader_subgroup_arithmetic : enable
#extension GL_KHR_shader_subgroup_ballot     : enable

// Sum across all subgroup lanes
float total = subgroupAdd(myValue);

// Predicate bitmask
uvec4 ballot = subgroupBallot(myValue > threshold);
uint  count  = subgroupBallotBitCount(ballot);

CUDA:

// Warp reduction
float sum = __shfl_down_sync(0xFFFFFFFF, val, 16);

// Predicate ballot
unsigned mask = __ballot_sync(0xFFFFFFFF, val > threshold);
int count     = __popc(mask);

2.3 Predicate / Filter Vectorization

Technology	Predicate API	Output
CPU AVX2	`_mm256_cmp_ps` + `_mm256_movemask_ps`	8-bit integer mask
CPU AVX-512	`_mm512_cmp_ps_mask`	16-bit `__mmask16`
Vulkan	`subgroupBallot(cond)`	`uvec4` bitmask
CUDA	`__ballot_sync(mask, pred)`	32-bit `unsigned int`

2.4 Reduction / Aggregation

Technology	Approach	Steps for 32 elements
CPU AVX2	`_mm256_hadd_ps` (horizontal)	~5 shuffles
Vulkan	`subgroupAdd(v)`	1 instruction
CUDA	`__shfl_down_sync` butterfly	5 steps, register-only

2.5 Graphics Pipeline Integration

Technology	Integration
CPU SIMD	None — explicit copy to GPU needed
Vulkan	Native — compute and graphics share the same pipeline. Vertex buffers can be produced by compute shaders directly.
CUDA	Interop only — via CUDA–Vulkan/OpenGL shared memory handles. Adds complexity and latency.

2.6 Decision Guide

Use case	Recommended
Small/hot data in CPU cache, latency-sensitive	CPU SIMD
Cross-platform GPU, mixed render+compute	Vulkan
Large-scale data-parallel analytics, ML, NVIDIA GPU	CUDA

3. Vectorization in Open-Source Databases — Examples

3.1 Apache Hive — ORC Vectorized Execution

Vectorized query execution processes a block of 1,024 rows at a time. Within the block, each column is stored as a vector — an array of a primitive data type.

-- Enable vectorized execution (required: ORC storage)
SET hive.vectorized.execution.enabled = true;
SET hive.vectorized.execution.reduce.enabled = true;

-- Verify: look for "Vectorized execution: true" in EXPLAIN output
EXPLAIN SELECT os, COUNT(*), AVG(duration)
FROM events
WHERE duration > 5000
GROUP BY os;

3.2 Apache Spark — Parquet Vectorized Reader

from pyspark.sql import SparkSession

spark = SparkSession.builder \
    .config("spark.sql.parquet.enableVectorizedReader", "true") \
    .config("spark.sql.parquet.columnarReaderBatchSize", "4096") \
    .config("spark.sql.columnVector.offheap.enabled", "true") \
    .getOrCreate()

events = spark.read.parquet("/tmp/events")

# Physical plan shows 'BatchScan' when vectorized reader is active
events.filter("duration > 5000") \
      .groupBy("os") \
      .agg({"duration": "avg"}) \
      .explain()

3.3 Apache Spark — Arrow-Backed Pandas UDFs

pandas_udf uses Apache Arrow to pass entire column batches between the JVM and Python, eliminating row-by-row serialization.

from pyspark.sql.functions import pandas_udf
from pyspark.sql.types import DoubleType
import pandas as pd

@pandas_udf(DoubleType())
def discount_score(duration: pd.Series) -> pd.Series:
    # Entire column as a Pandas Series — numpy vectorized op
    return (duration / duration.max()) * 100.0

events.withColumn("score", discount_score("duration")).show(5)

3.4 DuckDB — Vectorized Batch Execution with Apache Arrow

DuckDB's vectorized engine processes 1,024-row column chunks. It integrates with Apache Arrow datasets with filter pushdown before data leaves disk.

import duckdb
import pyarrow.dataset as ds

nyc = ds.dataset("nyc-taxi/", partitioning=["year", "month"])
con = duckdb.connect()

# Arrow filter pushdown + DuckDB vectorized engine
query = con.execute("""
    SELECT year, COUNT(*) AS trips, AVG(total_amount) AS avg_fare
    FROM nyc
    WHERE total_amount > 20 AND year > 2018
    GROUP BY year ORDER BY year
""")

# Stream results as Arrow RecordBatches
reader = query.fetch_record_batch()
for batch in reader:
    print(batch.to_pandas())

3.5 Vectorization Entry Points Summary

System	Key Config / API	Batch format	Requires
Hive	`hive.vectorized.execution.enabled=true`	1024-row column arrays	ORC / Parquet
Spark scan	`spark.sql.parquet.enableVectorizedReader=true`	`ColumnarBatch` (Arrow)	Parquet / ORC
Spark UDFs	`@pandas_udf`	Arrow `pd.Series` batch	`pyarrow`
DuckDB	Always-on	Arrow `RecordBatch`	Parquet / Arrow
DataFusion	Rust auto-SIMD	Arrow columnar + SIMD	Rust native

4. Vectorization API Usage Per Database

4.1 DuckDB

Layer	CUDA	CPU SIMD	Arrow	Vulkan/Metal
Scan / filter	None	Auto-vectorized (Clang, `-march=native`) SSE4.2 / AVX2 / AVX-512	Zero-copy via `duckdb_register_arrow()` / `fetch_record_batch()`	None
Aggregation	None	Implicit SIMD over 1024-row chunks	`rel.to_arrow_table()`	None

4.2 Apache Spark

Layer	CUDA	CPU SIMD	Arrow	Vulkan/Metal
Parquet scan	RAPIDS plugin (`com.nvidia.spark.SQLPlugin`)	`VectorizedParquetRecordReader` fills `ColumnarBatch`; JIT auto-vectorizes	`@pandas_udf` Arrow bridge	None
SQL / aggregation	cuDF operators (Sort, Join, GroupBy)	Whole-Stage Codegen (WSCG) fuses to tight JVM bytecode	Arrow UCX shuffle (RAPIDS)	None

4.3 Apache Hive

Layer	CUDA	CPU SIMD	Arrow	Vulkan/Metal
ORC/Parquet scan	None	JVM batch arrays (1024-row); JIT vectorizes `long[]`/`double[]` loops	None	None
Aggregation/join	None	LLAP pipeline; SSE/AVX via JIT	None	None

4.4 ClickHouse

Layer	CUDA	CPU SIMD	Arrow	Vulkan/Metal
Scan / filter	Not yet (GitHub #63392)	SSE4.2, AVX2, AVX-512 — runtime `cpuid` dispatch; 65,536-row blocks	Arrow Flight SQL (output only)	None
Aggregation	None	Explicit `_mm256_add_epi32` / `_mm512_add_epi64`; simdjson for JSON	None	None

4.5 Elasticsearch / OpenSearch

Layer	CUDA	CPU SIMD	Arrow	Vulkan/Metal
HNSW index build	NVIDIA cuVS (Elastic 8.18+) — 30–40% faster build	`simdvec` C++ library: float32, int8, bfloat16, BBQ; SSE/AVX-512/NEON	None	None
kNN distance	cuVS bulk dot-product / L2	Panama Vector API (`jdk.incubator.vector`)	None	None

4.6 Apache Arrow / DataFusion

Layer	CUDA	CPU SIMD	Arrow	Vulkan/Metal
Columnar execution	Auron (experimental, Spark plugin)	Rust LLVM auto-SIMD: AVX2 / AVX-512 / NEON; `std::simd` where explicit	Native `RecordBatch`; Arrow C Data Interface for JVM interop	None

Note: Vulkan and Metal are absent from all major open-source databases. These APIs lack the ecosystem libraries (cuBLAS, cuDF, CUB) that make CUDA practical for database workloads.

5. CUDA Image Processing APIs

5.1 NPP — NVIDIA Performance Primitives

NPP is the low-level GPU image & signal primitive library, part of the CUDA Toolkit.

Geometric Transforms

#include <nppi.h>

NppStreamContext ctx;
nppGetStreamContext(&ctx);  // CUDA < 13 only; see Section 6

// Resize (8-bit, 3-channel packed RGB)
nppiResize_8u_C3R_Ctx(
    pSrc, srcStep, srcSize,
    pDst, dstStep, dstSize,
    NPPI_INTER_LANCZOS,
    ctx);

// Rotate by angle (degrees) around (shiftX, shiftY)
nppiRotate_8u_C3R_Ctx(
    pSrc, srcSize, srcStep, srcROI,
    pDst,          dstStep, dstROI,
    45.0,           // angle
    0.0, 0.0,       // shiftX, shiftY
    NPPI_INTER_CUBIC,
    ctx);

// Affine warp — 2×3 matrix
double aCoeffs[2][3] = {{cos_a, -sin_a, tx},
                         {sin_a,  cos_a, ty}};
nppiWarpAffine_8u_C3R_Ctx(
    pSrc, srcSize, srcStep, srcROI,
    pDst,          dstStep, dstROI,
    aCoeffs, NPPI_INTER_LINEAR, ctx);

// Perspective warp — 3×3 homography
double hCoeffs[3][3] = {{ ... }};
nppiWarpPerspective_8u_C3R_Ctx(
    pSrc, srcSize, srcStep, srcROI,
    pDst,          dstStep, dstROI,
    hCoeffs, NPPI_INTER_LINEAR, ctx);

// Flip (horizontal / vertical / both)
nppiMirror_8u_C3R_Ctx(
    pSrc, srcStep,
    pDst, dstStep, roiSize,
    NPP_HORIZONTAL_AXIS, ctx);

Filters

// Gaussian blur
nppiFilterGauss_8u_C3R_Ctx(
    pSrc, srcStep, pDst, dstStep, roiSize,
    NPP_MASK_SIZE_5_X_5, ctx);

// Box (mean) blur
nppiFilterBox_8u_C3R_Ctx(
    pSrc, srcStep, pDst, dstStep, roiSize,
    {5, 5},   // oMaskSize
    {2, 2},   // oAnchor
    ctx);

// Median blur (requires scratch buffer)
int bufSize;
nppiFilterMedianGetBufferSize_8u_C1R(roiSize, {3,3}, &bufSize);
cudaMalloc(&pBuf, bufSize);
nppiFilterMedian_8u_C1R_Ctx(
    pSrc, srcStep, pDst, dstStep, roiSize,
    {3, 3}, {1, 1}, pBuf, ctx);

// Sobel edge detection
nppiFilterSobelHoriz_8u16s_C1R_Ctx(
    pSrc, srcStep, pDst, dstStep, roiSize,
    NPP_MASK_SIZE_3_X_3, ctx);

Morphology

// Dilation / erosion with user mask
NppiSize  maskSize = {3, 3};
NppiPoint anchor   = {1, 1};
Npp8u     pMask[]  = {1,1,1, 1,1,1, 1,1,1};  // 3×3 all-ones

nppiDilate_8u_C1R_Ctx(
    pSrc, srcStep, pDst, dstStep, roiSize,
    pMask, maskSize, anchor, ctx);

Color Conversion

// RGB → HSV
nppiRGBToHSV_8u_C3R_Ctx(pSrc, srcStep, pDst, dstStep, roiSize, ctx);

// YCbCr → RGB
nppiYCbCrToRGB_8u_C3R_Ctx(pSrc, srcStep, pDst, dstStep, roiSize, ctx);

// RGB → Grayscale (3-channel → 1-channel)
nppiRGBToGray_8u_C3C1R_Ctx(pSrc, srcStep, pDst, dstStep, roiSize, ctx);

// Video color formats
nppiNV12ToRGB_8u_P2C3R_Ctx(..., ctx);

NPP Stream Context Setup (CUDA 13+)

NppStreamContext ctx;
ctx.hStream              = myStream;
ctx.nCudaDeviceId        = deviceId;
ctx.nMultiProcessorCount = prop.multiProcessorCount;
ctx.nMaxThreadsPerMultiProcessor = prop.maxThreadsPerMultiProcessor;
ctx.nMaxThreadsPerBlock  = prop.maxThreadsPerBlock;
ctx.nSharedMemPerBlock   = prop.sharedMemPerBlock;
ctx.nCudaDevAttrComputeCapabilityMajor = prop.major;
ctx.nCudaDevAttrComputeCapabilityMinor = prop.minor;

5.2 CV-CUDA — Modern Batch-First GPU CV Library

CV-CUDA is NVIDIA's actively developed, recommended library for AI vision pre/post-processing. All operators work on NHWC tensors (Batch × Height × Width × Channels).

import cvcuda
import torch

# Create a batch of images as a NHWC CUDA tensor
batch = torch.zeros(8, 1080, 1920, 3, dtype=torch.uint8, device='cuda')
tensor = cvcuda.as_tensor(batch, "NHWC")

stream = cvcuda.Stream()
with stream:
    # Resize batch to 640×640
    resized = cvcuda.resize(tensor,
                            (8, 640, 640, 3),
                            cvcuda.Interp.LINEAR)

    # Rotate all images by 45 degrees
    rotated = cvcuda.rotate(resized,
                            angle_deg=45.0,
                            interp=cvcuda.Interp.CUBIC)

    # Gaussian blur
    blurred = cvcuda.gaussian(resized,
                              kernel_size=(5, 5),
                              sigma=(1.5, 1.5))

    # Bilateral filter (edge-preserving)
    filtered = cvcuda.bilateral_filter(resized,
                                       d=9,
                                       sigma_color=75,
                                       sigma_space=75)

    # Color convert BGR → RGB
    rgb = cvcuda.cvt_color(resized, cvcuda.ColorConversion.BGR2RGB)

    # Normalize for inference (subtract mean, divide by std)
    mean  = torch.tensor([0.485, 0.456, 0.406], device='cuda')
    std   = torch.tensor([0.229, 0.224, 0.225], device='cuda')
    normed = cvcuda.normalize(rgb,
                              base=mean,
                              scale=std,
                              flags=cvcuda.NormalizeFlags.SCALE_IS_STDDEV)

    # Warp affine (per-image transform matrices)
    matrices = torch.eye(2, 3, device='cuda').unsqueeze(0).repeat(8, 1, 1)
    warped = cvcuda.warp_affine(resized, matrices, cvcuda.Interp.LINEAR)

    # Warp perspective (homography)
    H = torch.eye(3, device='cuda').unsqueeze(0).repeat(8, 1, 1)
    proj = cvcuda.warp_perspective(resized, H, cvcuda.Interp.CUBIC)

    # Morphology
    kernel = cvcuda.StructuringElement(cvcuda.MorphShape.RECT, (3, 3))
    dilated = cvcuda.morphology(resized,
                                morphType=cvcuda.MorphType.DILATE,
                                structElem=kernel)

    # Remap (lens distortion / fisheye undistortion)
    map_x = torch.zeros(8, 1080, 1920, dtype=torch.float32, device='cuda')
    map_y = torch.zeros(8, 1080, 1920, dtype=torch.float32, device='cuda')
    remapped = cvcuda.remap(resized, map_x, map_y, cvcuda.Interp.LINEAR)

CV-CUDA Operator Reference

Operator	Python API	Notes
Resize	`cvcuda.resize()` / `cvcuda.hq_resize()`	HQResize: antialiasing, Lanczos, 3D tensor support
Rotate	`cvcuda.rotate()`	Cubic interp added in v0.10 (2× speedup)
Warp affine	`cvcuda.warp_affine()`	Per-image 2×3 matrix; supports inverse map
Warp perspective	`cvcuda.warp_perspective()`	Per-image 3×3 homography
Gaussian	`cvcuda.gaussian()`	Separable filter; σX, σY
Median	`cvcuda.median_blur()`	Non-linear noise removal
Bilateral	`cvcuda.bilateral_filter()`	Edge-preserving smooth
Color convert	`cvcuda.cvt_color()`	BGR↔RGB↔YUV↔NV12↔HSV…
Normalize	`cvcuda.normalize()`	Mean/std per-channel; DL preprocessing
Pad / border	`cvcuda.copy_make_border()`	Constant, reflect, replicate, wrap
Random crop	`cvcuda.random_resized_crop()`	Training augmentation
Morphology	`cvcuda.morphology()`	Dilate, Erode, Open, Close
Remap	`cvcuda.remap()`	Lens correction, fisheye, panoramic
SIFT	`cvcuda.sift()`	GPU feature detector
NMS	`cvcuda.nms()`	Non-maximum suppression

5.3 Image / Video Codecs

nvJPEG

#include <nvjpeg.h>

nvjpegHandle_t     handle;
nvjpegJpegState_t  state;

nvjpegCreate(NVJPEG_BACKEND_HARDWARE,  // HW engine on A100/H100
             nullptr, &handle);
nvjpegJpegStateCreate(handle, &state);

// Batched decode
nvjpegDecodeBatched(handle, state,
                    data_ptrs, lengths, batch_size,
                    destinations, params, stream);

NVDEC / NVENC (Video Codec SDK)

// Python wrapper (modern)
import PyNvVideoCodec as nvc

dec = nvc.CreateDecoder(gpuId=0, codec=nvc.CudaVideoCodec.H264)
enc = nvc.CreateEncoder(width, height, "h264", gpuId=0)

while True:
    frame = dec.DecodeSingleFrame()  # GPU tensor, no CPU round-trip
    # ... process frame with CV-CUDA ...
    enc.EncodeSingleFrame(frame)

5.4 Custom CUDA Kernels — Key Techniques

Texture Memory (sub-pixel interpolation)

// Create texture object for bilinear interpolation
cudaResourceDesc resDesc{};
resDesc.resType = cudaResourceTypePitch2D;
resDesc.res.pitch2D.devPtr = d_image;
resDesc.res.pitch2D.width  = width;
resDesc.res.pitch2D.height = height;
resDesc.res.pitch2D.pitchInBytes = pitch;
resDesc.res.pitch2D.desc = cudaCreateChannelDesc<float4>();

cudaTextureDesc texDesc{};
texDesc.filterMode      = cudaFilterModeLinear;  // HW bilinear
texDesc.addressMode[0]  = cudaAddressModeClamp;
texDesc.addressMode[1]  = cudaAddressModeClamp;
texDesc.normalizedCoords = 1;

cudaTextureObject_t tex;
cudaCreateTextureObject(&tex, &resDesc, &texDesc, nullptr);

// In kernel: hardware interpolation, no branches
__global__ void rotate_kernel(cudaTextureObject_t tex,
                               float4* out, float angle) {
    float u = ..., v = ...;  // compute rotated coordinates
    out[idx] = tex2D<float4>(tex, u, v);  // bilinear in hardware
}

Shared Memory Tile (convolution)

#define TILE_W 32
#define HALO   2   // kernel radius

__global__ void blur_kernel(const uchar3* src, uchar3* dst,
                             int W, int H) {
    __shared__ uchar3 tile[TILE_W + 2*HALO][TILE_W + 2*HALO];

    int tx = threadIdx.x, ty = threadIdx.y;
    int gx = blockIdx.x * TILE_W + tx - HALO;
    int gy = blockIdx.y * TILE_W + ty - HALO;

    // Load tile + halo into shared memory
    tile[ty][tx] = (gx >= 0 && gx < W && gy >= 0 && gy < H)
                   ? src[gy * W + gx] : uchar3{0,0,0};
    __syncthreads();

    // Apply filter using shared memory — no repeated global reads
    if (tx >= HALO && tx < TILE_W+HALO &&
        ty >= HALO && ty < TILE_W+HALO) {
        // ... box or Gaussian sum over tile[ty-HALO..ty+HALO][...]
    }
}

6. Legacy vs Modern CUDA Image APIs

6.1 What Changed in CUDA 13.0

All NPP functions without the _Ctx suffix were deprecated and hard-removed in CUDA 13.0. Additionally, nppGetStreamContext() and nppSetStream() were removed.

Old API (removed)	Replacement
`nppiResize_8u_C3R(...)`	`nppiResize_8u_C3R_Ctx(..., ctx)`
`nppiRotate_8u_C3R(...)`	`nppiRotate_8u_C3R_Ctx(..., ctx)`
`nppiFilterGauss_8u_C3R(...)`	`nppiFilterGauss_8u_C3R_Ctx(..., ctx)`
`nppGetStreamContext()`	Manually fill `NppStreamContext` struct
`nppSetStream(s)`	Pass `NppStreamContext` explicitly
NPP built-in JPEG (`nppicom`)	`nvJPEG` library

6.2 API Status Summary

API	Status	Notes
NPP non-`_Ctx` variants	Removed (CUDA 13)	All must be ported to `_Ctx` form
`nppGetStreamContext()`	Removed (CUDA 12.9/13)	Manually build `NppStreamContext`
`nppSetStream()`	Removed	Replaced by per-call context
NPP `_Ctx` variants	Current — valid	Supported in all current CUDA Toolkit
NPP+ (NPP_Plus)	Modern	C++ template API, multi-GPU, cleaner interface
CV-CUDA	Modern — recommended	Batch-first, PyTorch integration, active development
nvJPEG	Current	HW engine on Ampere+; use `NVJPEG_BACKEND_HARDWARE`
NVDEC / NVENC	Current	AV1 encode on Ada Lovelace/Hopper
`PyNvVideoCodec`	Modern	Replaces older VPF Python wrapper
DALI	Modern	Training-time augmentation + data loading
OpenCV CUDA (`cv::cuda::GpuMat`)	Superseded	Wraps NPP without batch support; use CV-CUDA for AI pipelines
CUDA texture for resize/rotate	Use with caution	Still valid; may be slower than dedicated kernels on Ampere+

6.3 Typical Modern AI Vision Pipeline

NVDEC (decode) → nvJPEG (JPEG decode)
    → cvcuda.resize() + cvcuda.normalize()
    → TensorRT inference
    → NPP color convert (if needed)
    → NVENC (re-encode)

All stages run asynchronously in separate CUDA streams, overlapping every phase simultaneously.

7. CPU–GPU Communication: Events and Messages

7.1 Fundamental Model

The CPU and GPU are separate processors connected by PCIe (or NVLink). They do not share a cache. All communication goes through:

Streams — command queues (CPU → GPU)
Events — timeline markers and signals (CPU ↔ GPU / stream ↔ stream)
Memory transfers — bulk data (CPU ↔ GPU)
In-kernel signaling — live notifications while a kernel is running
CUDA Graphs — pre-recorded pipeline replay

7.2 CUDA Streams

cudaStream_t stream;
cudaStreamCreate(&stream);

// Enqueue kernel — CPU returns IMMEDIATELY; GPU runs asynchronously
my_kernel<<<grid, block, 0, stream>>>(args);

// Block CPU until stream is empty
cudaStreamSynchronize(stream);

// Non-blocking poll — returns cudaSuccess or cudaErrorNotReady
cudaError_t status = cudaStreamQuery(stream);

// Block until ALL GPU work across all streams is done
cudaDeviceSynchronize();

cudaStreamDestroy(stream);

7.3 CUDA Events

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start, stream);
my_kernel<<<grid, block, 0, stream>>>(args);
cudaEventRecord(stop, stream);

// CPU blocks until 'stop' event fires
cudaEventSynchronize(stop);

// Elapsed time (ms, ~0.5 µs resolution)
float ms;
cudaEventElapsedTime(&ms, start, stop);

// Cross-stream dependency (GPU-side, no CPU involvement)
cudaEvent_t ev;
cudaEventCreateWithFlags(&ev, cudaEventDisableTiming);  // lower latency
cudaEventRecord(ev, streamA);           // mark end of producer
cudaStreamWaitEvent(streamB, ev, 0);    // consumer waits on GPU

cudaEventDestroy(start);
cudaEventDestroy(stop);
cudaEventDestroy(ev);

7.4 Async Memory Transfers (requires pinned memory)

// Allocate pinned (page-locked) host memory
float* h_data;
cudaMallocHost(&h_data, bytes);   // page-locked; enables true async

float* d_data;
cudaMalloc(&d_data, bytes);

// Async copy — returns immediately; copy runs on DMA engine
cudaMemcpyAsync(d_data, h_data, bytes, cudaMemcpyHostToDevice, stream);

// Overlap: kernel on stream1 while copy runs on stream0
my_kernel<<<grid, block, 0, stream1>>>(d_other);

cudaStreamSynchronize(stream0);  // wait for copy
cudaStreamSynchronize(stream1);  // wait for kernel

cudaFreeHost(h_data);

Important: Without pinned memory, cudaMemcpyAsync silently falls back to synchronous behavior.

7.5 Unified Memory

// Single pointer valid on both CPU and GPU
float* data;
cudaMallocManaged(&data, bytes);

// Prefetch to GPU before kernel (avoids page-fault stalls)
cudaMemPrefetchAsync(data, bytes, deviceId, stream);
my_kernel<<<grid, block, 0, stream>>>(data);

// Prefetch back to CPU
cudaMemPrefetchAsync(data, bytes, cudaCpuDeviceId, stream);
cudaStreamSynchronize(stream);
printf("Result: %f\n", data[0]);   // safe to read

// Memory usage hints
cudaMemAdvise(data, bytes, cudaMemAdviseSetReadMostly, deviceId);
cudaMemAdvise(data, bytes, cudaMemAdviseSetPreferredLocation, deviceId);
cudaMemAdvise(data, bytes, cudaMemAdviseSetAccessedBy, cudaCpuDeviceId);

cudaFree(data);

7.6 In-Kernel CPU↔GPU Signaling (live, while kernel runs)

GPU → CPU (progress notification)

// Host setup: zero-copy mapped memory
volatile unsigned long long *progress_h, *progress_d;
cudaHostAlloc((void**)&progress_h, sizeof(ull), cudaHostAllocMapped);
cudaHostGetDevicePointer((void**)&progress_d, (void*)progress_h, 0);

// GPU kernel: signal CPU after each block completes
__global__ void long_kernel(volatile unsigned long long* progress) {
    // ... do work ...

    if (threadIdx.x == 0) {
        atomicAdd((unsigned long long*)progress, 1);
        __threadfence_system();   // flush write through PCIe to CPU
    }
}

// Host: poll without blocking GPU
long_kernel<<<N_BLOCKS, 256>>>(progress_d);
while (*progress_h < N_BLOCKS) {
    printf("Progress: %llu / %d\n", *progress_h, N_BLOCKS);
    usleep(100000);  // 100 ms
}
cudaDeviceSynchronize();

__threadfence_system() is critical — without it the CPU may never see the atomic write due to GPU cache/reorder.

CPU → GPU (cancel / stop signal)

volatile int *stop_h, *stop_d;
cudaHostAlloc((void**)&stop_h, sizeof(int), cudaHostAllocMapped);
cudaHostGetDevicePointer((void**)&stop_d, (void*)stop_h, 0);
*stop_h = 0;

__global__ void persistent_kernel(volatile int* stop) {
    while (!(*stop)) {
        // ... process work items ...
    }
}

persistent_kernel<<<1, 32>>>(stop_d);

// From CPU thread or signal handler:
__sync_synchronize();
*stop_h = 1;        // GPU sees this on next poll (~1–5 µs PCIe latency)
cudaDeviceSynchronize();

`managed` variable (Pascal+)

__managed__ int flag = 0;   // accessible from both CPU and GPU

__global__ void signal_kernel() {
    // ... work ...
    atomicExch(&flag, 1);
    __threadfence_system();
}

signal_kernel<<<1, 32>>>();
while (flag == 0) {}   // CPU polls; driver manages coherence
cudaDeviceSynchronize();

7.7 CUDA Graphs — Eliminate Per-Launch Overhead

Each kernel launch from the CPU costs 5–20 µs in driver overhead. CUDA Graphs amortize this to a single call per replay.

// Capture a pipeline once
cudaGraph_t      graph;
cudaGraphExec_t  exec;
cudaStream_t     captureStream;
cudaStreamCreate(&captureStream);

cudaStreamBeginCapture(captureStream, cudaStreamCaptureModeGlobal);

// Record all operations
cudaMemcpyAsync(d_in, h_in, bytes, cudaMemcpyHostToDevice, captureStream);
kernel_a<<<grid, block, 0, captureStream>>>(d_in, d_mid);
kernel_b<<<grid, block, 0, captureStream>>>(d_mid, d_out);
cudaMemcpyAsync(h_out, d_out, bytes, cudaMemcpyDeviceToHost, captureStream);

cudaStreamEndCapture(captureStream, &graph);
cudaGraphInstantiate(&exec, graph, nullptr, nullptr, 0);

// Replay with ONE driver call per iteration
for (int i = 0; i < 10000; i++) {
    cudaGraphLaunch(exec, captureStream);
    cudaStreamSynchronize(captureStream);
}

cudaGraphExecDestroy(exec);
cudaGraphDestroy(graph);

8. CUDA Event Monitoring Internals: Polling vs Interrupts

There is no single "event monitoring thread" in CUDA by default. Whether polling, yielding, or interrupts are used depends on the scheduling mode set with cudaSetDeviceFlags.

8.1 Three Scheduling Modes

// Set BEFORE first CUDA API call (except cudaSetDevice)
cudaSetDeviceFlags(cudaDeviceScheduleSpin);         // busy-wait loop
cudaSetDeviceFlags(cudaDeviceScheduleYield);        // loop + sched_yield()
cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); // OS mutex + interrupt
cudaSetDeviceFlags(cudaDeviceScheduleAuto);         // driver heuristic (default)

Mode	CPU cost while waiting	Latency	Best for
`Spin`	100% of one core	Lowest (~µs)	Short GPU kernels < 1 ms, spare CPU core
`Yield`	Low (yields between polls)	Medium (~10s µs)	CPU has parallel work
`BlockingSync`	0% (thread sleeps)	Highest (~50–200 µs wake-up)	Long jobs > 10 ms, power-constrained
`Auto` (default)	Heuristic	Heuristic	Good default; override if needed

8.2 Layer 1 — Application: What the Modes Do

Spin: The CUDA runtime runs a tight busy-wait loop:

while (!GpuWorkFinished()) {}  // one CPU core at 100%

Yield: Same loop, but calls sched_yield() / SwitchToThread() each iteration so the OS can schedule other threads.

BlockingSync: The application thread blocks on a futex (Linux) / Win32 event. The CUDA driver has a dedicated background monitor thread per GPU context that wakes the application thread when the GPU fires an interrupt.

8.3 Layer 2 — Driver: Background Monitor Thread & ISR

Completion semaphore: When a kernel or DMA transfer finishes, the GPU Copy Engine writes a 32-bit completion value into a pinned memory page mapped into both GPU and CPU address spaces (the "channel semaphore"). Spin/yield modes poll this page directly with no system calls.

Driver monitor thread (BlockingSync only): The CUDA driver spawns one OS thread per GPU context. This thread blocks on an OS synchronization primitive. When the GPU fires an MSI interrupt, the nvidia.ko ISR runs and wakes this monitor thread, which then unblocks the application thread.

nvidia.ko ISR (interrupt handler): The NVIDIA Linux kernel module registers an ISR for the GPU's MSI-X interrupt line. The ISR runs atomically, reads the GPU interrupt status register via MMIO, and wakes the driver monitor thread. The ISR itself does minimal work — heavy lifting happens in the monitor thread.

nvidia-uvm.ko fault thread: A separate kernel module handles Unified Memory page faults. When a GPU thread accesses an unmapped page, the GPU MMU writes to a fault buffer and fires an interrupt. The UVM module ISR wakes a worker thread that migrates pages and replays faulting accesses, batching up to 32 faults per service cycle.

8.4 Layer 3 — Hardware: PCIe MSI-X Interrupts

PCIe MSI-X write (kernel completion): When the GPU Copy Engine finishes work, it writes a Message Signaled Interrupt — a small PCIe TLP write to a special CPU memory address pre-registered by the driver. This triggers a CPU APIC interrupt delivered to the core running the nvidia.ko ISR.

PCIe MSI-X write (UVM page fault): When a GPU SM causes a TLB miss, the GMMU writes a fault record to the replayable fault buffer and fires an MSI interrupt handled by nvidia-uvm.ko.

Semaphore write (no interrupt — Spin/Yield modes): The GPU CE writes a 32-bit value to the completion semaphore address. No interrupt is fired. The CPU polls this address. cudaEventQuery() is simply a memory read of this location — zero system call overhead.

8.5 End-to-End Flow (BlockingSync Path)

CPU app thread
  calls cudaEventSynchronize(e)
  → blocks on OS futex/mutex
        ↓
  CUDA driver thread
    sleeps on futex, waiting for ISR to wake it
        ↓
  GPU Copy Engine
    finishes kernel, writes completion semaphore
    fires PCIe MSI-X write transaction
        ↓
  PCIe bus
    delivers interrupt write to CPU APIC
        ↓
  nvidia.ko ISR (runs atomically on a CPU core)
    reads GPU interrupt status register (MMIO)
    wakes CUDA driver monitor thread
        ↓
  CUDA driver monitor thread
    reads completion semaphore value
    signals the futex/mutex
        ↓
  CPU app thread
    unblocked from cudaEventSynchronize()
    resumes execution

Spin path (short circuit):

GPU writes completion semaphore value
    ↓
CPU spin loop reads that address
  while (*semaphore != expected) {}   ← polling this memory address
    ↓
Loop exits — no ISR, no driver thread, no system call

8.6 `cudaEventQuery()` vs `cudaEventSynchronize()`

// Non-blocking poll — just a memory read, zero system call
cudaError_t s = cudaEventQuery(event);
if (s == cudaSuccess)         puts("done");
if (s == cudaErrorNotReady)   puts("still running");

// Blocking wait — mode-dependent (spin / yield / OS mutex)
cudaEventSynchronize(event);  // returns when event has fired

Use cudaEventQuery() in a loop when the CPU has other useful work to do between checks. Use cudaEventSynchronize() when you simply need to wait.

Appendix: Quick Reference

Scheduling Mode Selection

Situation	Recommended mode
GPU kernels < 1 ms, latency critical	`cudaDeviceScheduleSpin`
CPU runs parallel threads while GPU works	`cudaDeviceScheduleYield`
GPU kernels > 10 ms (training, long analytics)	`cudaDeviceScheduleBlockingSync`
Embedded / Tegra / power-constrained	`cudaDeviceScheduleBlockingSync`
Default / don't know	`cudaDeviceScheduleAuto`

Key Include Headers

#include <cuda_runtime.h>   // cudaStream, cudaEvent, cudaMemcpy, etc.
#include <cub/cub.cuh>      // CUB device-wide primitives
#include <thrust/...>       // Thrust algorithms
#include <nppi.h>           // NPP image processing
#include <nvjpeg.h>         // JPEG encode/decode
#include <cufft.h>          // 2D/3D FFT

GPU vs CPU Vectorization Quick Lookup

Task	CPU	GPU (CUDA)
SIMD add 8 floats	`_mm256_add_ps`	`float4` load + scalar add per thread
Predicate filter	`_mm256_cmp_ps` + `movemask`	`__ballot_sync`
Horizontal sum	`_mm256_hadd_ps` (multi-step)	`__shfl_down_sync` (5 steps)
Batch scan	`thrust::copy_if` on CPU	`cub::DeviceSelect::If`
Sort	`std::sort`	`cub::DeviceRadixSort`

Generated from a full technical session on GPU vectorization, image processing, and CPU–GPU communication.

Table of Contents​

1. CUDA APIs for Vectorized Query Processing​

1.1 Vectorized Memory Access — Vector Types​

1.2 CUB (CUDA UnBound)​

1.3 Thrust​

1.4 Warp Shuffle — Warp-Level Reduction​

1.5 Warp Ballot — Predicate Vectorization​

1.6 SWAR (SIMD Within A Register) + 128-bit Loads​

1.7 CUDA Streams + Async Memory Copy​

1.8 cuFFT — Frequency-Domain Filtering​

1.9 API Summary Table​

2. CPU SIMD vs Vulkan vs CUDA — Comparison​

2.1 Execution Model​

2.2 Key Vectorization APIs​

2.3 Predicate / Filter Vectorization​

2.4 Reduction / Aggregation​

2.5 Graphics Pipeline Integration​

2.6 Decision Guide​

3. Vectorization in Open-Source Databases — Examples​

3.1 Apache Hive — ORC Vectorized Execution​

3.2 Apache Spark — Parquet Vectorized Reader​

3.3 Apache Spark — Arrow-Backed Pandas UDFs​

3.4 DuckDB — Vectorized Batch Execution with Apache Arrow​

3.5 Vectorization Entry Points Summary​

4. Vectorization API Usage Per Database​

4.1 DuckDB​

4.2 Apache Spark​

4.3 Apache Hive​

4.4 ClickHouse​

4.5 Elasticsearch / OpenSearch​

4.6 Apache Arrow / DataFusion​

5. CUDA Image Processing APIs​

5.1 NPP — NVIDIA Performance Primitives​

Geometric Transforms​

Filters​

Morphology​

Color Conversion​

NPP Stream Context Setup (CUDA 13+)​

5.2 CV-CUDA — Modern Batch-First GPU CV Library​

CV-CUDA Operator Reference​

5.3 Image / Video Codecs​

nvJPEG​

NVDEC / NVENC (Video Codec SDK)​

5.4 Custom CUDA Kernels — Key Techniques​

Texture Memory (sub-pixel interpolation)​

Shared Memory Tile (convolution)​

6. Legacy vs Modern CUDA Image APIs​

6.1 What Changed in CUDA 13.0​

6.2 API Status Summary​

6.3 Typical Modern AI Vision Pipeline​

7. CPU–GPU Communication: Events and Messages​

7.1 Fundamental Model​

7.2 CUDA Streams​

7.3 CUDA Events​

7.4 Async Memory Transfers (requires pinned memory)​

7.5 Unified Memory​

7.6 In-Kernel CPU↔GPU Signaling (live, while kernel runs)​

GPU → CPU (progress notification)​

CPU → GPU (cancel / stop signal)​

__managed__ variable (Pascal+)​

7.7 CUDA Graphs — Eliminate Per-Launch Overhead​

8. CUDA Event Monitoring Internals: Polling vs Interrupts​

8.1 Three Scheduling Modes​

8.2 Layer 1 — Application: What the Modes Do​

8.3 Layer 2 — Driver: Background Monitor Thread & ISR​

8.4 Layer 3 — Hardware: PCIe MSI-X Interrupts​

8.5 End-to-End Flow (BlockingSync Path)​

8.6 cudaEventQuery() vs cudaEventSynchronize()​

Appendix: Quick Reference​

Scheduling Mode Selection​

Key Include Headers​

GPU vs CPU Vectorization Quick Lookup​

Table of Contents