CUDA Standard Algorithms » Parallel Transforms

Transform a Range of Items

Parallel-transform algorithm applies the given transform function to a range of items and store the result in another range specified by two iterators, first and last. The task created by tf::cuda_transform(P&& p, I first, I last, O output, C op) represents a parallel execution for the following loop:

while (first != last) {
  *output++ = op(*first++);
}

The following example creates a transform kernel that transforms an input range of N items to an output range by multiplying each item by 10.

tf::cuda_transform(
  input, input + N, output, [] __device__ (int x) { return x*10; }
);

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.

Transform Two Ranges of Items

You can transform two ranges of items to an output range through a binary operator. The task created by tf::cuda_transform(P&& p, I1 first1, I1 last1, I2 first2, O output, C op) represents a parallel execution for the following loop:

while (first1 != last1) {
  *output++ = op(*first1++, *first2++);
}

The following example creates a transform kernel that transforms two input ranges of N items to an output range by summing each pair of items in the input ranges.

// output[i] = input1[i] + inpu2[i]
tf::cuda_transform(
  input1, input1+N, input2, output, []__device__(int a, int b) { return a+b; }
);

Invoke Parallel-Transform Kernels Asynchronously

By default, tf::cuda_transform blocks until all kernels finish. You can use tf::cuda_transform_async to invoke parallel-transform kernels asynchronously and explicitly synchronize them at another place of your program.

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

for(size_t i=0; i<1000; i++) input[i] = i;

// Launch transform kernel asynchronously
tf::cuda_transform_async(
  p, input, input+1000, output [] __device__ (int item) { return 10*item; }
);
p.synchronize();  // now output[i] = input[i] * 10