cudaMallocとcudaMallocPitchの秘密を教えます。

 からの許可を得て転載しています。 http://blog.csdn.net/bendanban/article/details/7646306

タイトルにある2つの関数のテストに興味があったのですが、アライメントアクセスデータを満たすために、通常はcudamallocPitchを使うことが多く、その方が高い効率をもたらすと考えています。あ、ここに自分のマシンで実行できるテストプログラムがありますが、この2つの関数が同じ場合もあることがわかります。

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime_api.h>

int main(int argc, char **argv)
{
	// device pointers.
	float *d_pitch;
	float *d_normal;

	// matrix size.
	size_t cols = 63;
	size_t rows = 16;
	
	size_t pitch = 0;
	
	// alloc the data form gpu memory.
	cudaMallocPitch((void**)&d_pitch, &pitch, cols*sizeof(float), rows);
	cudaMalloc((void**)(&d_normal), rows*cols*sizeof(float));
	
	// test the data address.
	fprintf(stdout, "row size(in bytes) = %.2f*128.\n", pitch/128.0f);
	fprintf(stdout, "the head address of d_pitch mod 128 = %x.\n", ((unsigned int)d_pitch)%128);
	fprintf(stdout, "the head address of d_normal mod 128 = %x.\n", ((unsigned int)d_normal)%128);
	
	cudaFree(d_pitch);
	cudaFree(d_normal);

	getchar();
	return 0;
}

上記のプログラムは、以下のように実行されます。

row size(in bytes) = 28.00*128.
the head address of d_pitch mod 128 = 0.
the head address of d_normal mod 128 = 0. 

私は何度も実験していますが、上記の実験結果から、実験のパラメータをどう変えても、2つのメモリ要求関数が返すデータの先頭アドレスは128と256の整数倍であり、GPU上の各計算ユニットのデータは一度に2の累乗でグローバルにロードでき、これらのデータロードのアドレスも2の累乗でなければならないと推測します。 そこでwarpはグローバルメモリ内のデータを使う際にアライメントの原則に従ってデータをロードするようにして、より効率を上げる必要があると思います。アライメントの原則は、CUDAのプログラミングマニュアルに記載されています。