私は、Tesla M2090 (Fermi) に基づくクラスターと、K20Xm (Kepler) に基づく別のクラスターを使用しています。Fermi クラスターで起動した私のカーネルは、Kepler よりも 2.5 倍高速です。このカーネルは、キー -arch=sm_35 --ptxas-options=-v を使用して Kepler クラスター用にコンパイルされました。結果は次のとおりです。
ptxas info : Compiling entry function '_Z22_repack_one_thread_8_2ILb1EEviPtPPh' for 'sm_35'
ptxas info : Function properties for _Z22_repack_one_thread_8_2ILb1EEviPtPPh
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 18 registers, 344 bytes cmem[0]
したがって、スレッドごとに 18 個のレジスタと 0 バイトの共有メモリを使用して 1024 個のスレッドを使用すると、マルチプロセッサの占有率が 100% になります。
Kepler に基づくノードのパフォーマンスが遅くなる理由として考えられるのは何ですか?
ありがとうございました。
ヴォイツェフ
アップデート
私のカーネル
template <bool nocheck>
__global__ void _repack_one_thread_8_2 (int size, word *input, byte **outputs)
{
int idx = blockDim.x * blockIdx.x + threadIdx.x;
if (nocheck || idx * 8 < size)
{
word *ptr = input + idx * 4;
byte bytes[8] = {0,0,0,0,0,0,0,0};
int i, j;
for (i = 0; i < 4; i++, ptr++)
{
word b = *ptr;
for (j = 0; j < 8; j++)
bytes[j] |= (((b >> (j * 2)) & 3) << (i * 2));
}
for (i = 0; i < 8; i++)
outputs[i][idx] = bytes[i];
}
}
Kepler のコンパイル コマンド
nvcc -arch=sm_35 --ptxas-options=-v -c -O3 -I.. -o
Fermi のコンパイル コマンド
nvcc -arch=sm_20 --ptxas-options=-v -c -O3 -I.. -o