【初心者向け】CUDAの__shfl関数の使い方

warp（ワープ）の組み込み関数の__shflの使用方法について説明します。
シャッフル関数を使用する前に、warpについて理解する必要があるので、ワープから順に説明していきます。

Warp(ワープ)について
__shfl_sync関数

Warp(ワープ)について

スレッドはWarpと呼ばれる32スレッドの集合で動きます。
例えば、グリッドサイズ=1、ブロックサイズ=96でカーネルを呼び出した時、以下のように3つのワープが動きます。

CUDAのカーネルを高速に動かすには、warpがとても重要です。
warpには以下の特徴があります。

warp内のスレッドは同時に動く
warp内のスレッドは連続しているアドレスにアクセスするとき最も高速になる
warp内のスレッドはスレッド間でデータのやり取りが可能

今回は③の特徴を使ったプログラムになります。

__shfl_sync関数

データのやり取りをする__shfl_sync関数の説明です。
似たような関数に__shfl関数というものがありますが、最新でより安全な__shfl_sync関数を使いましょう。
__shfl_sync関数は公式サイトによると以下のように使用します。

実際に動かしてみます。


#include <iostream>
#include <vector>
#include <stdio.h>
#include <numeric>
#include <windows.h>
#include <cuda.h>
#include <algorithm>

#pragma comment(lib, "winmm.lib")




__global__
void kernel(float* d_val, float* d_out)
{

    unsigned int tid = blockIdx.x * blockDim.x+threadIdx.x;


    float val = d_val[tid];

    val = __shfl_sync(0xffffffff,val, tid + 1, warpSize);

    d_out[tid] = val;
}



int main()
{


    int n = 32;
    float* d_val;
    float* d_out;
    float* h_out;
    std::vector<float> h_val;



    for (int i = 0; i < n; i++) h_val.push_back(float(i+1));

    cudaMalloc(&d_val, sizeof(float) * n);
    cudaMalloc(&d_out, sizeof(float) * n);
    cudaMemcpy(d_val, h_val.data(), sizeof(float) * n, cudaMemcpyHostToDevice);

    h_out = (float*)malloc(sizeof(float));

    int block = 32;
    int gird =1;
    

    kernel << < gird, block >> > (d_val, d_out);

    cudaMemcpy(h_out, d_out, sizeof(float)*n, cudaMemcpyDeviceToHost);

    for (int i = 0; i < n; i++)
    {
        printf("%f\n", h_out[i]);
    }




}

0xffffffffはおまじないだと思ってください。

出力結果

__shfl_sync関数は以下のようにアクセスしているのが分かります。