bulk

We have a lot parallel work to do, but all we have are these puny threads. Let's Bulk up!

Bulk lets you describe a parallel task as a hierarchical grouping of execution agents. Individually, these agents are like tiny, lightweight threads, but when grouped together pack some serious muscle.

We can launch parallel groups of agents with bulk::async:

hello_world.cu:

#include <iostream>
#include <cstdio>
#include <bulk/bulk.hpp>

struct hello
{
  __host__ __device__
  void operator()()
  {
    printf("Hello world!\n");
  }

  __host__ __device__
  void operator()(bulk::parallel_group<> &g)
  {
    printf("Hello world from thread %d\n", g.this_exec.index());
  }
};

int main()
{
  // just launch one agent to say hello
  bulk::async(bulk::par(1), hello());

  // launch 32 agents in parallel
  // bulk::root stands in for the root of the agent hierarchy
  // the hello functor uses this to identify each agent within its group
  bulk::async(bulk::par(32), hello(), bulk::root);

  cudaDeviceSynchronize();

  return 0;
}

From here it's a trivial exercise to get to SAXPY:

#include <iostream>
#include <cassert>
#include <bulk/bulk.hpp>
#include <thrust/device_vector.h>

struct saxpy
{
  __host__ __device__
  void operator()(bulk::agent<> &self, float a, float *x, float *y)
  {
    int i = self.index();
    y[i] = a * x[i] + y[i];
  }
};

int main()
{
  size_t n = 1 << 24;
  thrust::device_vector<float> x(n, 1);
  thrust::device_vector<float> y(n, 1);

  float a = 13;

  // pass bulk::root.this_exec so the saxpy functor receives
  // the current execution agent directly
  bulk::async(bulk::par(n), saxpy(), bulk::root.this_exec, a, thrust::raw_pointer_cast(x.data()), thrust::raw_pointer_cast(y.data()));

  assert(y == thrust::device_vector<float>(n, 14));

  std::cout << "Nice SAXPY. Do you work out?" << std::endl;

  return 0;
}

Algorithms built with Bulk are fast.

reduce Performance

inclusive_scan Performance

merge Performance

reduce_by_key Performance