Implementing the Down-Sweep Phase in Exclusive Scan

Introduction

In our last lesson, we constructed the first half of the work-efficient Blelloch scan algorithm: the up-sweep phase. We transformed an array of input data into a conceptual binary tree of partial sums within shared memory. After that phase, the last element of our array holds the total sum of all elements, and the other elements store the intermediate sums from the reduction.

We're now ready to complete the algorithm. This lesson focuses on the down-sweep phase, which takes the reduction tree we just built and elegantly traverses it in the opposite direction—from the root back to the leaves—to produce the final exclusive prefix sum. By the end of this lesson, you will have a complete, work-efficient, single-block scan kernel, a truly fundamental tool in high-performance computing.

The Down-Sweep Phase: From Partial Sums to Final Scan

The up-sweep phase left us with a shared memory array full of useful partial sums. The goal of the down-sweep is to use these sums to compute the final exclusive scan, where each output element y_i is the sum of all input elements before index i.

The process works in $lo g_{2} N$ steps, mirroring the up-sweep but moving top-down. It starts by "seeding" the root of the tree (the last element of our array) with the identity element of the operation, which is 0 for addition. Then, at each level, a parent node passes information down to its two children. The logic is as follows:

The parent node knows the prefix sum for the entire block of data it represents.
It passes this prefix sum directly to its left child.
It passes the sum of its own prefix sum plus the partial sum of the left child (which was computed during the up-sweep) to its right child.

This propagation ensures that every element eventually receives the sum of all elements that preceded it in the original array.

Let's look at a diagram of the full algorithm. We covered the "up-sweep" portion last time; now our focus is on the "down-sweep" portion.

The complete Blelloch scan for an 8-element array. The top half is the up-sweep (reduction), and the bottom half is the down-sweep. Notice how the down-sweep starts by zeroing out the last element and works with decreasing offsets.

To see this logic in action, a dynamic walkthrough is invaluable.

Blelloch Scan - Intro to Parallel Programming

This video from Udacity provides a clear, animated explanation of the down-sweep phase, including the special operator and a concrete numerical example.

Watch from the end of the example. Pay close attention to these three key points: How the rightmost element is initialized to the identity (0 for sum). The unique "downsweep operator" that takes two inputs and produces two outputs, which corresponds to the left/right child logic. How intermediate results from the up-sweep phase are "dragged down" and used in combination with values computed during the down-sweep itself.

The Down-Sweep Algorithm and In-Place Implementation

The "pass to left child, pass sum to right child" logic can be implemented very efficiently in-place using a clever swap-and-add pattern. To understand this, let's look at the formal pseudocode.

Chapter 39. Parallel Prefix Sum (Scan) with CUDA

The "GPU Gems 3" chapter we reviewed in the last lesson also provides excellent pseudocode and a diagram for the down-sweep phase.

Focus on the section that begins right after the up-sweep description. Study Figure 39-4 and the pseudocode in "Example 4". This pseudocode shows a swap-and-add pattern. Try to relate t = x[k + 2^d - 1] and the subsequent lines to the conceptual "pass to left child, pass sum to right child" logic we discussed.

As the pseudocode shows, at each step d of the down-sweep, for each pair of elements ai and bi being operated on, the algorithm performs:

t = temp[ai] (Save the value of the left child, which is a partial sum from the up-sweep).
temp[ai] = temp[bi] (The left child receives the value from the parent, which is temp[bi]).
temp[bi] = temp[bi] + t (The right child receives the value from the parent plus the original value of the left child).

This sequence perfectly executes the down-sweep logic in place within the shared memory array.

Translating the Down-Sweep to a CUDA Kernel

Now we can write the CUDA code for the down-sweep, which will follow the up-sweep phase within our kernel.

First, after the up-sweep is complete and all threads are synchronized, a single thread sets the last element of the shared memory array to zero. This initializes the root of the down-sweep.

The down-sweep loop then begins. Unlike the up-sweep loop which counted down (d >>= 1), the down-sweep loop counts up (d *= 2), as the number of active threads doubles at each level moving away from the root. The offset is now divided by two at each step to select the correct elements.

Here is the code segment for the down-sweep, which would appear directly after the up-sweep code from our previous lesson.

// temp[] is our __shared__ float array of size n
// 'offset' is carried over from the up-sweep, its final value is n/2.

// After the up-sweep, clear the last element (the root of the tree).
// This is done by one thread, as it's a single write operation.
if (threadIdx.x == 0) {
    temp[n - 1] = 0;
}

// The down-sweep loop traverses "down" the tree.
// 'd' represents the number of active thread pairs at each level.
for (int d = 1; d < n; d *= 2)
{
    // The offset shrinks as we move down the tree levels.
    offset >>= 1;
    __syncthreads(); // Sync to ensure temp[n-1]=0 is visible and previous level is complete.

    // 'd' threads are active in this step.
    if (threadIdx.x < d)
    {
        // Identify the indices for the left and right children.
        int ai = offset * (2 * threadIdx.x + 1) - 1;
        int bi = offset * (2 * threadIdx.x + 2) - 1;

        // The core in-place swap-and-add operation.
        float t = temp[ai];   // Save the left child's partial sum.
        temp[ai] = temp[bi];  // Left child gets the exclusive prefix from the parent.
        temp[bi] += t;        // Right child gets the parent's prefix + left child's sum.
    }
}
__syncthreads(); // Ensure the final down-sweep writes are visible to all threads.

// At this point, the temp[] array contains the complete exclusive scan result.
// The final step is to write temp[] back to global memory.

By tracing the indices ai and bi with the changing offset, you can see how this code systematically performs the swap-and-add operations shown in the diagrams, propagating the correct prefix sums throughout the array until the final exclusive scan is computed.

The Complete Single-Block Scan Kernel

By combining the code from the previous lesson (data loading and up-sweep) with the code from this lesson (down-sweep and data storing), we arrive at a complete, work-efficient, single-block exclusive scan.

Here is the full kernel structure for reference. It is based on the implementation from NVIDIA's "GPU Gems 3" (LINK, section 2), which you should now be able to understand completely.

__global__ void prescan(float *g_odata, const float *g_idata, int n)
{
    // Allocate shared memory dynamically based on block size
    extern __shared__ float temp[];

    int thid = threadIdx.x;

    // Load two elements per thread from global to shared memory
    int i = 2 * thid;
    if (i < n) temp[i] = g_idata[i];
    if (i + 1 < n) temp[i + 1] = g_idata[i + 1];

    // --- UP-SWEEP PHASE ---
    int offset = 1;
    for (int d = n >> 1; d > 0; d >>= 1)
    {
        __syncthreads();
        if (thid < d)
        {
            int ai = offset * (2 * thid + 1) - 1;
            int bi = offset * (2 * thid + 2) - 1;
            temp[bi] += temp[ai];
        }
        offset *= 2;
    }

    // --- DOWN-SWEEP PHASE ---
    if (thid == 0) { temp[n - 1] = 0; }

    for (int d = 1; d < n; d *= 2)
    {
        offset >>= 1;
        __syncthreads();
        if (thid < d)
        {
            int ai = offset * (2 * thid + 1) - 1;
            int bi = offset * (2 * thid + 2) - 1;
            float t = temp[ai];
            temp[ai] = temp[bi];
            temp[bi] += t;
        }
    }
    __syncthreads();

    // Write the results from shared memory back to global memory
    if (i < n) g_odata[i] = temp[i];
    if (i + 1 < n) g_odata[i + 1] = temp[i + 1];
}

Conclusion

In this lesson, you completed the implementation of one of the most important parallel primitives. By adding the down-sweep phase to the up-sweep, we have a complete and highly efficient Blelloch scan algorithm.

Key Takeaways:

Down-Sweep Purpose: To distribute the partial sums computed during the up-sweep to generate the final exclusive scan result.
Core Logic: The algorithm traverses the conceptual reduction tree from root to leaves. A parent node passes its value (an exclusive prefix sum) to its left child, and passes the sum of its value and the left child's up-sweep sum to its right child.
Initialization: The down-sweep is seeded by setting the last element of the array (the root of the tree) to the operator's identity element (0 for addition).
In-Place Implementation: This complex propagation can be implemented efficiently in-place with a swap-and-add pattern: t = a; a = b; b += t;.
Complete Algorithm: The combination of the up-sweep and down-sweep, with __syncthreads() at each level traversal and between the two phases, forms the complete, work-efficient scan algorithm.

We have now added reduction, histogramming, and scan to our toolkit of parallel primitives. In the next lesson, we will explore two more fundamental patterns, gather and scatter, which are essential for reorganizing data in parallel according to arbitrary mappings.