Texture Memory使いたかったから勘弁してください

この記事書いてる途中にOpera for Ubuntu11.10が2度クラッシュしてデータ消えたのでもうUbuntuで記事書いてるときにOperaは使わないと心に決めました。

もう書くのめんどくさくなってきたのでさくっと書きます。

CUDAにはTexture MemoryがあってTexture Unitからのみアクセス可能ですが、Texture Unitには専用にキャッシュ領域があるので、アクセスが集中するようなデータをそこへ格納することでキャッシュの恩恵を受けてアクセスの高速化を図れるのでは、ということです。

例えば、行列ベクトル積 Ax = b においてベクトル x は計算時に複数回参照されることが分かります。そこにキャッシュの恩恵を適用すれば、アクセスにかかる時間が短くなり、計算が早くなるのではないか、というお話のようです。

データの正規化とか線形補間とかできるとかなんとかかかれていますが、私がお邪魔している研究室的には巨大な一次連立方程式を解くのが主目的なようなので、逆に補完されると困るって話なようなのでそういうのは使いません。

Texture Memoryを使用するには、textureクラステンプレートを宣言します。*1

  texture<float, 1, cudaReadModeElementType> float_tex;

パラメータが、テクスチャ内の数値型(int, int2, floatなど), 次元数, 読み取りモードの指定(cudaReadModeElementType, cudaReadModeNormalizedFloat)の順に指定します。

cudaReadModeNormalizedFloatが正規化を行うタイプで、私が使う予定はないです(多分...)

cudaReadModeElementTypeを指定しておくと、配列と似た感じのアクセスが可能になります。整数型のインデックスを指定したい場合、tex1Dfetch関数を使います。*2

次元数は、まあ1次元でベクトル、2次元で行列、とか考えておくと分かるのではないでしょうか。

後はテクスチャにバインドしたりなんだったりなんですが、面倒なのでコード書けばいいですね。

static const std::size_t N = ...;

texture<float, cudaTextureType1D, cudaReadModeElementType> texture_;

template<class T>
__device__
void kernel_(int n, T const* x, T* y) {
  int i = ...;
  y[i] = x[i] * tex1Dfetch(texture_, i);
}

float *dev_vec; // デバイスのメモリ
float *dev_x, *dev_y;

cudaBindTexture(0, texture_, dev_vec,
                cudaCreateChannelDesc<float>(), sizeof(float) * N);
kernel_<<<..., ...>>>(N, dev_x, dev_y);
cudaUnbindTexture(texture_);

floatとかだったらこれで良いのですが、doubleはそのままでは渡せないようです。

このページのProgramming questions 23. "Can I read double precision floats from texture?"にTexture Memoryでのdoubleの使い方が書かれています。CUDA 2.1での解答なんですが、多分今もこの対応でいいのでしょう。

doubleを使おうとするとこうなります。

static const std::size_t N = ...;

texture<int2, cudaTextureType1D, cudaReadModeElementType> texture_;

template<class T>
__device__
void kernel_(int n, T const* x, T* y) {
  int i = ...;
  int2 v = tex1Dfetch(texture_, i);
  y[i] = x[i] * __hiloint2double(v.y, v.x);
}

double *dev_vec; // デバイスのメモリ
double *dev_x, *dev_y;

cudaBindTexture(0, texture_, dev_vec,
                cudaCreateChannelDesc<int2>(), sizeof(int2) * N);
kernel_<<<..., ...>>>(N, dev_x, dev_y);
cudaUnbindTexture(texture_);

アクセスめんどうですね。しかも型によって切り替えるし。研究室の人は多倍長整数なども実装したいと言っているらしく、多倍長にしたらまたアクセス方法が変わってきますね。

計算したい型情報から、テクスチャメモリの1要素のデータ型と、実際の計算で使うデータ型、後tex1Dfetchを使ってアクセスしてくれるようなTraitsクラスを書いた。

// デフォルト, テクスチャのデータ型と計算時のデータ型が同一
template<class DataType>
struct traits_ {
  typedef DataType value_type;
  typedef DataType texture_value_type;

  static const int fetch_dim = cudaTextureType1D;

  __device__
  value_type operator()(
    texture<texture_value_type, fetch_dim, cudaReadModeElementType>& vt,
    int index
  ) const {
    return tex1Dfetch(vt, index);
  }
};

// double 用
template<>
struct traits_<double> {
  typedef double value_type;
  typedef int2   texture_value_type;

  static const int fetch_dim = cudaTextureType1D;

  __device__
  value_type operator()(
    texture<texture_value_type, fetch_dim, cudaReadModeElementType>& vt,
    int index
  ) const {
    int2 v = tex1Dfetch(vt, index);
    return __hiloint2double(v.y, v.x);
  }
};

typedef double                    using_float_type;
typedef traits_<using_float_type> texture_traits;

texture<
    texture_traits::texture_float_type
  , texture_traits::fetch_dim
  , cudaReadModeElementType
> texture_;

static const std::size_t N = ...;

using_float_type *dev_vec;
using_float_type *dev_x, *dev_y;

cudaBindTexture( 0, texture_, dev_vec
               , cudaCreateChannelDesc<texture_traits::texture_value_type>()
               , sizeof(texture_traits::value_type) * N );
kernel_<<<..., ...>>>(N, dev_x, dev_y);
cudaUnbindTexture(texture_);

独自に定義した多倍長整数なんかでも、traitsクラスを特殊化すれば対応できると思う。

説明が結構端折ったり、足りてない部分があると思ったけど、CUDA C Programming Guide 4.1とか、ちゃんとドキュメント読まないと死ぬって事がよくわかってきた昨今です。

*1:CUDA C Programming Guideなどを読んでいると多分クラステンプレートという解釈で良いと思われます

*2:tex1DLayerdとかあるんだけどこれはなんだろうか