CUDA Scatter/Gather Operations Guide

December 25, 2025 9 min By RightNow AI Team

Introduction

Scatter and gather operations move data between positions based on index tensors. Gather reads from indexed positions (like embedding lookup), while scatter writes to indexed positions (like sparse gradient updates). Graph neural networks heavily use these for message passing.

This guide covers efficient implementations for both operations with focus on avoiding race conditions in scatter.

Common Performance Issues

• Race conditions when multiple threads scatter to same index
• Atomic operations serialize access, hurting performance
• Unsorted indices cause poor cache utilization

Optimization Techniques

1. Segment-based Scatter

Sort indices, process segments without atomics.

2. Deterministic Scatter

Use sorting for reproducible results.

3. Vectorized Gather

Use float4 for consecutive dimension access.

Implementation Comparison

Before (Naive Implementation)

Atomics are correct but serialize when indices collide.

// Scatter with atomics - correct but slow
__global__ void scatter_add_atomic(float* src, int* indices, float* dst,
                                   int n, int dim) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= n) return;

int target = indices[idx];
    for (int d = 0; d < dim; d++) {
        atomicAdd(&dst[target * dim + d], src[idx * dim + d]);
    }
}

After (Optimized Implementation)

Sorting enables segment-based reduction without atomics.

// Sort indices, then process segments
// Step 1: Sort (indices, values) by indices
// Step 2: Find segment boundaries
// Step 3: Reduce each segment

#include <cub/cub.cuh>

void scatter_add_sorted(float* src, int* indices, float* dst,
                        int n, int dim, void* temp) {
    // Sort indices and get permutation
    int* sorted_indices;
    int* permutation;
    cub::DeviceRadixSort::SortPairs(temp, temp_bytes,
                                     indices, sorted_indices,
                                     permutation, n);

// Now consecutive threads access consecutive memory
    // Segment reduce without atomics
    scatter_segment_reduce<<<grid, block>>> (
        src, sorted_indices, permutation, dst, n, dim);
}

__global__ void scatter_segment_reduce(float* src, int* sorted_idx,
                                       int* perm, float* dst, int n, int dim) {
    // Identify segment boundaries
    // Reduce within segment using warp shuffle
    // Single thread per segment writes result
}

Performance Results

Frequently Asked Questions

When to use atomic vs sorted scatter?

Atomic is simpler for low collision rates (<5% duplicates). Sorted is faster for high collision (embeddings, GNNs) but has O(n log n) sort overhead.