どこにも言及されていない (少なくとも私が見ることができる) のは、インライン CUDA カーネルに公開されているライブラリ関数です。
具体的には、個別に GPU にオフロードする価値のない小さな/愚かな行列乗算を行っていますが、この乗算を含むアルゴリズムのより大きなセクションをオフロードしています。自分の linalg 関数を使うのが好きな人はいませんでした。
TLDR PyCUDA でインライン カーネルを使用しているときに、どのライブラリを使用できますか?
どこにも言及されていない (少なくとも私が見ることができる) のは、インライン CUDA カーネルに公開されているライブラリ関数です。
具体的には、個別に GPU にオフロードする価値のない小さな/愚かな行列乗算を行っていますが、この乗算を含むアルゴリズムのより大きなセクションをオフロードしています。自分の linalg 関数を使うのが好きな人はいませんでした。
TLDR PyCUDA でインライン カーネルを使用しているときに、どのライブラリを使用できますか?
私は何も知りません、そして私はいつもそれがあれば便利だと思っていました.
私が通常扱う問題 (有限要素法で発生する小さな行列とテンソル) のサイズについては、操作を行う C++ テンプレートを書きました。関数をテンプレート化することで、コンパイラはコンパイル時にトリップ カウントを知ることができ、ループをアンロールして、結果または中間結果をレジスタに保持できます。これは、カーネル スループットにとって非常に効率的である傾向があります。したがって、行列-行列積は次のように宣言されます
template < typename Real, unsigned int l, unsigned int m, unsigned int n >
__device__ __host__
void matmul(const Real *a,
const Real *b,
Real *c)
{
for(int i=0; i<l; i++) {
for(int j=0; j<n; j++) {
Real dotprod = Real(0);
for(int k=0; k<m; k++) {
dotprod += a[idx2c(i,k,l)] * b[idx2c(k,j,m)];
}
c[idx2c(i,j,l)] = dotprod;
}
}
}
私のカーネルで発生する種類のサイズ (2x2、3x3、4x4、8x8、9x9) については、上記を実行してコンパイルを機能させることは、私が試した他のアプローチと同じくらい良いようです。スレッド レベルでは CUDA は事実上スカラーであるため、これらの種類の小さな操作を高速化するために使用できるベクトル プリミティブやそのようなものはありません。