CUDA Standard Algorithms » Parallel Scan

Scan a Range of Items

tf::cuda_inclusive_scan computes an inclusive prefix sum operation using the given binary operator over a range of elements specified by [first, last). The term "inclusive" means that the i-th input element is included in the i-th sum. The following code computes the inclusive prefix sum over an input array and stores the result in an output array.

const size_t N = 1000000;
int* input  = tf::cuda_malloc_shared<int>(N);  // input  vector
int* output = tf::cuda_malloc_shared<int>(N);  // output vector

// initializes the data
for(size_t i=0; i<N; i++) 
  input[i] = rand();
} 
// computes inclusive scan over input and stores the result in output
tf::cuda_inclusive_scan(tf::cudaDefaultExecutionPolicy{}, 
  input, input + N, output, [] __device__ (int a, int b) { return a + b; }
);
// verifies the result
for(size_t i=1; i<N; i++) {
  assert(output[i] == output[i-1] + input[i]);
}

On the other hand, tf::cuda_exclusive_scan computes an exclusive prefix sum operation. The term "exclusive" means that the i-th input element is NOT included in the i-th sum.

// computes exclusive scan over input and stores the result in output
tf::cuda_exclusive_scan(tf::cudaDefaultExecutionPolicy{}, 
  input, input + N, output, [] __device__ (int a, int b) { return a + b; }
);
// verifies the result
for(size_t i=1; i<N; i++) {
  assert(output[i] == output[i-1] + input[i-1]);
}

Scan a Range of Transformed Items

tf::cuda_transform_inclusive_scan transforms each item in the range [first, last) and computes an inclusive prefix sum over these transformed items. The following code multiplies each item by 10 and then compute the inclusive prefix sum over 1000000 transformed items.

const size_t N = 1000000;
int* input  = tf::cuda_malloc_shared<int>(N);  // input  vector
int* output = tf::cuda_malloc_shared<int>(N);  // output vector

// initializes the data
for(size_t i=0; i<N; i++) 
  input[i] = rand();
} 
// computes inclusive scan over transformed input and stores the result in output
tf::cuda_transform_inclusive_scan(tf::cudaDefaultExecutionPolicy{}, 
  input, input + N, output, 
  [] __device__ (int a, int b) { return a + b; },  // binary scan operator
  [] __device__ (int a) { return a*10; }           // unary transform operator
);
// verifies the result
for(size_t i=1; i<N; i++) {
  assert(output[i] == output[i-1] + input[i] * 10);
}

Similarly, tf::cuda_transform_exclusive_scan performs an exclusive prefix sum over a range of transformed items. The following code computes the exclusive prefix sum over 1000000 transformed items each multipled by 10.

const size_t N = 1000000;
int* input  = tf::cuda_malloc_shared<int>(N);  // input  vector
int* output = tf::cuda_malloc_shared<int>(N);  // output vector

// initializes the data
for(size_t i=0; i<N; i++) 
  input[i] = rand();
} 

// computes exclusive scan over transformed input and stores the result in output
tf::cuda_transform_exclusive_scan(tf::cudaDefaultExecutionPolicy{}, 
  input, input + N, output, 
  [] __device__ (int a, int b) { return a + b; },  // binary scan operator
  [] __device__ (int a) { return a*10; }           // unary transform operator
);

// verifies the result
for(size_t i=1; i<N; i++) {
  assert(output[i] == output[i-1] + input[i-1] * 10);
}

Invoke Scan Kernels Asynchronously

By default, scan functions block until all kernels finish. You can invoke these kernels asynchronously and explicitly synchronize them at another place of your program. Since our scan kernels rely on additional device memory, you need to provide a buffer of size in bytes equal to (or larger than) the value returned by tf::cuda_scan_buffer_size.

const size_t N = 1000000;
int* input  = tf::cuda_malloc_shared<int>(N);  // input  vector
int* output = tf::cuda_malloc_shared<int>(N);  // output vector

// initializes the data
for(size_t i=0; i<N; i++) 
  input[i] = rand();
} 

// queries the required buffer size to scan N elements using the given policy
auto bytes = tf::cuda_scan_buffer_size<tf::cudaDefaultExecutionPolicy, int>(N);
void* buffer = tf::cuda_malloc_device<std::byte>(bytes);

// invoke the inclusive scan kernels asynchronously (same for other scan variants)
tf::cuda_inclusive_scan_async(tf::cudaDefaultExecutionPolicy{my_stream},
  input, input+N, output, [] __device__ (int a, int b) { return a + b; }, buffer
);
// here, output may not be ready yet, as kernels are still running ...

// synchronize the stream to make output ready
cudaStreamSynchronize(my_stream);

The allocated buffer must remain alive until the asynchronous call completes.