CUDAのベクトル型

ベクトル型

CUDAではfloatやdoubleをいくつかまとめたベクトル型が提供されています． 言ってしまえばただの構造体のなのですが，アラインメントあたりが工夫されていたり等，特に理由がない限りオレオレ構造体ではなくこちらを使ったほうが嬉しいかもしれません．

例えばこんな感じ． (include/vector_types.h)

struct __device_builtin__ __builtin_align__(16) float4
{
     float x, y, z, w;
};

ベクトル型の種類

• char1 char2 char3 char4
• uchar1 uchar2 uchar3 uchar4
• short1 short2 short3 short4
• ushort1 ushort2 ushort3 ushort4
• int1 int2 int3 int4
• uint1 uint2 uint3 uint4
• long1 long2 long3 long4
• ulong1 ulong2 ulong3 ulong4
• longlong1 longlong2 longlong3 longlong4
• ulonglong1 ulonglong2 ulonglong3 ulonglong4
• float1 float2 float3 float4
• double1 double2 double3 double4

たくさんありますね．

詳しいアラインメント等はCUDAのDocument[1]を読んでみてください.

実は複素数型であるcuComplexなどもこのベクトル型がtypedefされたものです．

使い方

ただの構造体なので要素x,y,z,wにアクセスすれば普通に使えます．

他にも

const auto f4 = make_float4(0.0f, 1.0f, 2.0f, 3.0f);

などのベクトルを作る関数も用意されています．

重要なことですが，half2とは異なりSIMD命令は用意されていないので計算自体は速くなりません．

ベクトル型の利点

グローバルメモリの読み書きが高速化されるかもしれない

float 4つのオレオレ構造体だと読み込み部分のPTXはld.global 4つになるのに対し，公式のベクトル型では ld.global.v4.f32というベクトル読み込み命令が発行されます．(v2, v4のみ)

__ldg関数が利用可能

一部のベクトル型[2]ではRead-Onlyキャッシュ経由でグローバルメモリにアクセスする__ldg関数に対応しています． 読み込みが速くなるかもしれません．

実験

float<1,2,3,4>, double<1,2,3,4>の読み込み速度比較

• 各型の変数をグローバルメモリから各スレッドが1つ読み込む
• gitlab.momo86.net : mutsuki/cuda-vector
• GTX 1080

double3とdouble4の読み込みがほぼ同じ時間なことやfloat4の方がfloat3より速く読み込まれていることは興味深い

終わりに

グローバルメモリアクセスはボトルネックになりやすいのでうまくベクトル型を使いこなせると幸せな気分になれそうですね．

参考サイト

1. CUDA C Programming Guide - Built-in Vector Types
2. CUDA C Programming Guide - Read-Only Data Cache Load Function