CUDAでMonte Carlo Simulation -「Cooperative Groups」

Cooperative Groupsについて学んだのでメモ。

詳細は以下参照。

Cooperative Groups: Flexible CUDA Thread Programming | NVIDIA Developer Blog
Cooperative Groups: Flexible CUDA Thread Programming | NVIDIA Developer Blog

 
普通にCUDAでスレッド並列しようとするとブロック単位で動かすことになってしまうけど、cooperative_group.hを使うことで、より粒度を細かくすることができるという代物。
BlockとThreadの間に一つ階層が増えるイメージだと思ってる。

以下はCUDAのサンプルにあるMonteCarlo SimulationのReductionのサンプルコード。
サンプルコードでは２つの配列を別々にReductionしてるが、１つのOption Priceを求めたいときは１つの配列で十分。

・this_thread_block()でBlock内のスレッドを扱うインスタンスを召喚
・tiled_partition<32>()でスレッドを32個ごとに分割する
・thread_rank()でthreadのidを指定
・sync()でスレッドを同期

・下記コードで用いられているReduction手法はインタリーブペア方式と呼ばれるもので、配列を半分に分けて分割した１要素目同士、２要素目同士…を足し合わせていく。２ループ目では配列をさらに半分に分割して同様に足し合わせていく。そして最終的に合計値が求まる。

#include <cooperative_groups.h>

namespace cg = cooperative_groups;

template<class T, int SUM_N, int blockSize>
__device__ void sumReduce(T *sum, T *sum2, cg::thread_block &cta, cg::thread_block_tile<32> &tile32, __TOptionValue *d_CallValue)
{
const int VEC = 32; //Threadのサイズ
const int tid = cta.thread_rank();

T beta = sum[tid];
T beta2 = sum2[tid];
T temp, temp2;

for (int i = VEC/2; i > 0; i>>=1)  //sum[tid]とsum[tid+i]を足し合わせる。１ループ目でVEC/2回足し合わせ、２ループ目でVEC/2/2回足し合わせていく
{
if (tile32.thread_rank() < i)
{
temp = sum[tid+i];
temp2 = sum2[tid+i];
beta += temp;
beta2 += temp2;
sum[tid] = beta;
sum2[tid] = beta2;
}
cg::sync(tile32); //分割したPartition Threadを同期
}
cg::sync(cta); //Block内のThreadを同期

if (tid == 0) //最終結果
{
beta = 0;
beta2 = 0;
for (int i = 0; i < blockDim.x; i += VEC)  //Partition間で合計値を計算
{
beta += sum[i];
beta2 += sum2[i];
}
__TOptionValue t = {beta, beta2};
*d_CallValue = t;
}
cg::sync(cta);
}