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]を読んでみてください.

使い方

ただの構造体なので要素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関数に対応しています． 読み込みが速くなるかもしれません．

グローバルメモリへの結合アクセスが可能

グローバルメモリにあるfloat1, float2, float3, float4型の領域を各スレッドが読み込むと，読み込みは float1, float2, float4, float3の順に速いことが多いです． float3よりfloat4が速いのは結合アクセスが行われているためだと考えられます．

なので，例えば物理シミュレーション等で3次元の速度ベクトルを保持しようと思ったらfloat3よりfloat4を使ったほうがメモリアクセスは良くなるかもしれないですね．(余った1要素には違う物理量を入れるといいかも)
でもでも，読み込み時と書き込み時で高速化への程度が異なるので，ちゃんと実験をしてから使用するかどうか決めたほうがいいかもしれないです．

実験

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

終わりに

参考サイト

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