CUDA マシンで発生しているインデックスの問題をデバッグしようとしています
Cuda Machine Info:
{1->{Name->Tesla C2050,クロック レート->1147000,計算能力->2.,GPU オーバーラップ->1,最大ブロック ディメンション->{1024,1024,64},最大グリッド ディメンション->{65535 ,65535,65535},ブロックあたりの最大スレッド数->1024,ブロックあたりの最大共有メモリ->49152,合計定数メモリ->65536,ワープ サイズ->32,最大ピッチ->2147483647,ブロックあたりの最大レジスタ数->32768,テクスチャ アライメント ->512、マルチプロセッサ カウント ->14、コア カウント ->448、実行タイムアウト ->0、統合 ->False、ホスト メモリをマップ可能 ->True、コンピューティング モード ->デフォルト、Texture1D 幅 ->65536、Texture2D幅->65536、Texture2D 高さ->65535、Texture3D 幅->2048、Texture3D 高さ->2048、Texture3D 深さ->2048、Texture2D 配列幅->16384、Texture2D 配列高さ->16384、Texture2D 配列スライス->2048、 Surface Alignment->512,Concurrent Kernels->True,ECC有効->True,合計メモリ->2817982462},
このコードは、CUDA が使用しているインデックスに等しい 3D 配列の値を設定するだけです。
__global __ void cudaMatExp(
float *matrix1, float *matrixStore, int lengthx, int lengthy, int lengthz){
long UniqueBlockIndex = blockIdx.y * gridDim.x + blockIdx.x;
long index = UniqueBlockIndex * blockDim.z * blockDim.y * blockDim.x +
threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x +
threadIdx.x;
if (index < lengthx*lengthy*lengthz) {
matrixStore[index] = index;
}
}
何らかの理由で、3D 配列の次元が大きくなりすぎると、インデックス作成が停止します。
さまざまなブロックサイズを試しました (blockDim.x by blockDim.y by blockDim.z):
8x8x8 は、配列次元 12x12x12 までの正しいインデックス付けのみを提供します
9x9x9 は、配列次元 14x14x14 までの正しいインデックス付けのみを提供します
10x10x10 は、配列次元 15x15x15 までの正しいインデックス付けのみを提供します
これらよりも大きな次元の場合、さまざまなブロック サイズのすべてが最終的に再び増加し始めますが、dim^3-1 の値に達することはありません (これは、cuda スレッドが到達する最大インデックスです)。
以下に、この動作を示すプロットをいくつか示します。
例: これは、x 軸に 3D 配列の次元 (x x x) をプロットし、y 軸に cuda 実行中に処理される最大インデックス番号をプロットします。この特定のプロットは、ブロックのサイズが 10x10x10 の場合です。
そのプロットを生成する (Mathematica) コードを次に示しますが、これを実行したときは、1024x1x1 のブロック次元を使用しました。
CUDAExp = CUDAFunctionLoad[codeexp, "cudaMatExp",
{{"Float", _,"Input"}, {"Float", _,"Output"},
_Integer, _Integer, _Integer},
{1024, 1, 1}]; (*These last three numbers are the block dimensions*)
max = 100; (* the maximum dimension of the 3D array *)
hold = Table[1, {i, 1, max}];
compare = Table[i^3, {i, 1, max}];
Do[
dim = ii;
AA = CUDAMemoryLoad[ConstantArray[1.0, {dim, dim, dim}], Real,
"TargetPrecision" -> "Single"];
BB = CUDAMemoryLoad[ConstantArray[1.0, {dim, dim, dim}], Real,
"TargetPrecision" -> "Single"];
hold[[ii]] = Max[Flatten[
CUDAMemoryGet[CUDAExp[AA, BB, dim, dim, dim][[1]]]]];
, {ii, 1, max}]
ListLinePlot[{compare, Flatten[hold]}, PlotRange -> All]
これは同じプロットですが、x^3 をプロットして、あるべき場所と比較しています。配列の次元が 32 を超えると発散することに注意してください。
3D 配列の次元をテストし、インデックスがどこまで進んでいるかを調べ、dim^3-1 と比較します。たとえば、dim=32 の場合、cuda の最大インデックスは 32767 (32^3 -1) ですが、dim=33 の場合、35936 (33^3 -1) になるはずの cuda 出力が 33791 になります。33791-32767 = 1024 = blockDim.x であることに注意してください。
質問：
Mathematica のブロック次元よりも大きい次元の配列に正しくインデックスを付ける方法はありますか?
さて、ビット乗算のエラーを防ぐためにインデックス方程式で __mul24(threadIdx.y,blockDim.x) を使用する人がいることを知っていますが、私の場合は役に立たないようです。
また、デフォルトでは計算機能 1.0 用にコンパイルされているため、コードを -arch=sm_11 でコンパイルする必要があると誰かが言及しているのを見ました。これが Mathematica の場合かどうかはわかりません。CUDAFunctionLoad[] は 2.0 機能でコンパイルすることを知っていると思います。誰でも知っていますか？
どんな提案も非常に役に立ちます！