CUDA-メモリ割り当て

CUDAはGPUの並列計算ユニットに依存するソフトウェアとハードウェアからなる並列計算システムと考えられ,CUDAにはクラスCのAPIがあり,プログラム作成が容易である.CPUとGPUの異機種系に依存し、CPU上で環境初期化、メモリ割り当て、データ転送をシリアル実行し、GPU上で並列計算を実行する.
メモリ割当て
1、1次元

int *dev_ans = 0;
cudaMalloc((void**)&dev_ans, d.y * sizeof(int));

パラメータ1:メモリにおいて開かれた空間のポインタ(用語:GPUデバイス側データポインタ)
パラメータ2:スペースサイズ、バイト単位
2、2 D

int *dev_mat = 0;
int pitch;
cudaMallocPitch((void**)&dev_mat, (size_t *)&pitch, d.x * sizeof(int), d.y);

パラメータ1:GPUデバイス側データポインタ
パラメータ2:1行のデータの実空間サイズ(バイト)【このパラメータは戻り値の取得】であり、GPUにおいて256バイトのアドレス(address=0256512…)から連続的にアクセスすることが最も効率的であるため、各行の実割り当てのサイズは割り当てるサイズより大きい
パラメータ3:各行に割り当てるスペースサイズ
パラメータ4:行列行数
 
メモリコピー
1、1次元

cudaMemcpy(ans, dev_ans, d.y * sizeof(int), cudaMemcpyDeviceToHost);

パラメータ1:ターゲットデータアドレス
パラメータ2:ソースデータアドレス
パラメータ3:データサイズ
パラメータ4:コピータイプ(ホストからホスト、ホストからデバイス、デバイスからホスト、デバイスからデバイス)
2、2 D

cudaMemcpy2D(dev_mat, pitch, mat, d.x*sizeof(int), d.x*sizeof(int), d.y, cudaMemcpyHostToDevice);

パラメータ1:ターゲットデータアドレス
パラメータ2:pitch、割り当て空間の行幅(バイト単位)
パラメータ3:ソースデータアドレス
パラメータ4:pitch,割当て空間の行幅(バイト単位)
パラメータ5:データをコピーする必要がある実際の行幅(バイト単位)
パラメータ6:データの行数(バイト単位じゃないよ!)
パラメータ7:データコピータイプ
注意:pitchは線形記憶空間の行幅であり、データの行幅ではなく、デバイス側でpitchがデータの行幅以上であり、ホスト側でpitch==データの行幅である.
 
メモリアクセス
ホスト内のメモリアクセスがc++のアクセスであることは言うまでもありませんが、ビデオメモリのアクセス方法(kernelでのアクセス)を見てみましょう.

__global__ void addKernel(int *mat, int *ans, size_t pitch)
{
    int bid = blockIdx.x;
    int tid = threadIdx.x;
    __shared__ int data[8];
    int *row = (int*)((char*)mat + bid*pitch);
    data[tid] = row[tid];
    __syncthreads();
    for (int i = 4; i > 0; i /= 2) {
        if (tid < i)
            data[tid] = data[tid] + data[tid + i];
        __syncthreads();
    }
    if (tid == 0)
        ans[bid] = data[0];
}

1次元:
ans[index]ダイレクトアクセス
2 D:
アクセス先行の初期アドレスint*row=(int*)((char*)mat+bid*pitch)を先に計算します.
次に、この行の対応する要素row[index]にアクセスします.
 
メモリの解放

cudaFree(dev_mat)