CUDA Standard Algorithms » Parallel Iterations

Index-based Parallel Iterations

Index-based parallel-for performs parallel iterations over a range [first, last) with the given step size. The task created by tf::cuda_for_each_index represents a kernel of parallel execution for the following loop:

// positive step: first, first+step, first+2*step, ...
for(auto i=first; i<last; i+=step) {
  callable(i);
}
// negative step: first, first-step, first-2*step, ...
for(auto i=first; i>last; i+=step) {
  callable(i);
}

Each iteration i is independent of each other and is assigned one kernel thread to run the callable. Since the callable runs on GPU, it must be declared with a __device__ specifier. The following example creates a kernel that assigns each entry of data to 1 over the range [0, 100) with step size 1.

auto data = tf::cuda_malloc_shared<int>(100);

// assigns each element in data to 1 over the range [0, 100) with step size 1
tf::cuda_for_each_index(tf::cudaDefaultExecutionPolicy{},
  0, 100, 1, [data] __device__ (int idx) { data[idx] = 1; }
);

Iterator-based Parallel Iterations

Iterator-based parallel-for performs parallel iterations over a range specified by two STL-styled iterators, first and last. The task created by tf::cuda_for_each represents a parallel execution of the following loop:

for(auto i=first; i<last; i++) {
  callable(*i);
}

The two iterators, first and last, are typically two raw pointers to the first element and the next to the last element in the range in GPU memory space. The following example creates a for_each kernel that assigns each element in gpu_data to 1 over the range [data, data + 1000).

auto data = tf::cuda_malloc_shared<int>(1000);

// assigns each element in data to 1 over the range [0, 1000) with step size 1
tf::cuda_for_each(tf::cudaDefaultExecutionPolicy{},
  data, data + 1000, [] __device__ (int& item) { item = 1; }
);

Each iteration is independent of each other and is assigned one kernel thread to run the callable. Since the callable runs on GPU, it must be declared with a __device__ specifier.

Invoke Parallel-Iteration Kernels Asynchronously

By default, tf::cuda_for_each and tf::cuda_for_each_index blocks until all kernels finish. You can use tf::cuda_for_each_async and tf::cuda_for_each_index_async to invoke parallel-iteration kernels asynchronously and explicitly synchronize them at another place of your program.

auto p = tf::cudaDefaultExecutionPolicy(my_stream);
auto data = tf::cuda_malloc_shared<int>(1000);

// Launch for_each kernel asynchronously
tf::cuda_for_each_async(
  p, data, data + 1000, [] __device__ (int& item) { item = 1; }
);
p.synchronize();  // now all items in data are 1

// Launch for_each_index kernel asynchronously
tf::cuda_for_each_index_async(
  p, 0, 1000, 1, [data] __device__ (int idx) { data[idx] = -1; }
);
p.synchronize();  // now all items in data are -1