CUDA Standard Algorithms » Parallel Reduction

Reduce a Range of Items with an Initial Value

tf::cuda_reduce performs a parallel reduction over a range of elements specified by [first, last) using the binary operator bop and stores the reduced result in result. It represents the parallel execution of the following reduction loop on a GPU:

while (first != last) {
  *result = bop(*result, *first++);
}

The variable result participates in the reduction loop and must be initialized with an initial value. The following code performs a parallel reduction to sum all the numbers in the given range with an initial value 1000:

const size_t N = 1000000;
int* res = tf::cuda_malloc_shared<int>(1);  // result
int* vec = tf::cuda_malloc_shared<int>(N);  // vector

// initializes the data
*res = 1000;
for(size_t i=0; i<N; i++) 
  vec[i] = i;
} 

// *res = 1000 + (0 + 1 + 2 + 3 + 4 + ... + N-1)
tf::cuda_reduce(tf::cudaDefaultExecutionPolicy{},
  vec, vec + N, res, [] __device__ (int a, int b) { return a + b; }
);

Reduce a Range of Items without an Initial Value

tf::cuda_uninitialized_reduce performs a parallel reduction over a range of items without an initial value. This method represents a parallel execution of the following reduction loop on a GPU:

*result = *first++;  // no initial values to participate in the reduction loop
while (first != last) {
  *result = bop(*result, *first++);
}

The variable result is directly assigned the reduced value without any initial value participating in the reduction loop. The following code performs a parallel reduction to sum all the numbers in the given range without any initial value:

const size_t N = 1000000;
int* res = tf::cuda_malloc_shared<int>(1);  // result
int* vec = tf::cuda_malloc_shared<int>(N);  // vector

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

// *res = 0 + 1 + 2 + 3 + 4 + ... + N-1
tf::cuda_uninitialized_reduce(tf::cudaDefaultExecutionPolicy{},
  vec, vec + N, res, [] __device__ (int a, int b) { return a + b; }
);

Reduce a Range of Transformed Items with an Initial Value

tf::cuda_transform_reduce performs a parallel reduction over a range of transformed elements specified by [first, last) using a binary reduce operator bop and a unary transform operator uop. It represents the parallel execution of the following reduction loop on a GPU:

while (first != last) {
  *result = bop(*result, uop(*first++));
}

The variable result participates in the reduction loop and must be initialized with an initial value. The following code performs a parallel reduction to sum all the transformed numbers multiplied by 10 in the given range with an initial value 1000:

const size_t N = 1000000;
int* res = tf::cuda_malloc_shared<int>(1);  // result
int* vec = tf::cuda_malloc_shared<int>(N);  // vector

// initializes the data
*res = 1000;
for(size_t i=0; i<N; i++) 
  vec[i] = i;
} 

// *res = 1000 + (0*10 + 1*10 + 2*10 + 3*10 + 4*10 + ... + (N-1)*10)
tf::cuda_reduce(tf::cudaDefaultExecutionPolicy{},
  vec, vec + N, res, 
  [] __device__ (int a, int b) { return a + b; },
  [] __device__ (int a) { return a*10; }
);

Reduce a Range of Transformed Items without an Initial Value

tf::cuda_transform_uninitialized_reduce performs a parallel reduction over a range of transformed items without an initial value. This method represents a parallel execution of the following reduction loop on a GPU:

*result = *first++;  // no initial values to participate in the reduction loop
while (first != last) {
  *result = bop(*result, uop(*first++));
}

The variable result is directly assigned the reduced value without any initial value participating in the reduction loop. The following code performs a parallel reduction to sum all the transformed numbers multiplied by 10 in the given range without any initial value:

const size_t N = 1000000;
int* res = tf::cuda_malloc_shared<int>(1);  // result
int* vec = tf::cuda_malloc_shared<int>(N);  // vector

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

// *res = 0*10 + 1*10 + 2*10 + 3*10 + 4*10 + ... + (N-1)*10
tf::cuda_uninitialized_reduce(tf::cudaDefaultExecutionPolicy{},
  vec, vec + N, res, 
  [] __device__ (int a, int b) { return a + b; },
  [] __device__ (int a) { return a*10; }
);

Invoke Reduction Kernels Asynchronously

By default, tf::cuda_reduce and other variants blocks until all kernels finish. You can use tf::cuda_reduce_async to invoke these kernels asynchronously and explicitly synchronize them at another place of your program. Since our reduction 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_reduce_buffer_size.

const size_t N = 1000000;
int* res = tf::cuda_malloc_shared<int>(1);  // result
int* vec = tf::cuda_malloc_shared<int>(N);  // vector

// initializes the data
for(size_t i=0; i<N; i++) 
  vec[i] = i;
} 
*res = 0;

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

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

// synchronize the stream
cudaStreamSynchronize(my_stream);

assert(*res == (N-1)*N/2);

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