CUDAメモリ型メモリ

CUDAのメモリタイプ。

各スレッドは独自のレジスタとローカルメモリを持ちます。

各スレッドブロックには共有メモリが1つずつあります。

すべてのスレッドがグローバルメモリにアクセスできる。

また、すべてのスレッドがアクセス可能な読み取り専用メモリ：定数メモリ、テクスチャメモリ

1. レジスター登録

　　レジスタはGPU上のキャッシュです。基本単位はレジスタファイルで、各レジスタファイルのサイズは32ビットです。

　　カーネルのローカル変数（単純型）は、レジスターに割り当てるのが第一の選択です。

　　特徴：スレッドごとのプライベート、高速。

2.ローカルメモリ ローカルメモリ

　　レジスタを使い切ると、データはローカルメモリに格納されます。各スレッドで使用するレジスタが多すぎる場合、大きな構造体や配列が宣言されている場合、コンパイラが配列のサイズを決定できない場合、スレッドのプライベートデータはローカルメモリに割り当てられます。

　　特徴：スレッドごとのプライベート、キャッシュがない、遅い。

　　注意：ローカル変数を宣言する場合は、レジスタに代入可能な変数にするよう心がけましょう。

　　unsigned int mt[3];

　　unsigned int mt0, mt1, mt2 に置き換えてください。

3.共有メモリ 共有メモリ

　　同一ブロック内の全スレッドで読み出し、書き込みが可能

　　特徴：ブロック内の全スレッドに共通、共有メモリへのアクセスはレジスタとほぼ同等に高速。

[cpp］ 表示プレーン コピー

印刷 ?

//u(i)= u(i)^2 + u(i-1)  
//Static  
__global__ example(float* u) {  
    int i=threadIdx.x;  
    __shared__ int tmp[4];  
     tmp[i]=u[i];  
     u[i]=tmp[i]*tmp[i]+tmp[3-i];  
}  
  
int main() {  
    float hostU[4] = {1, 2, 3, 4};  
    float* devU;  
    size_t size = sizeof(float)*4;  
    cudaMalloc(&devU, size);  
    cudaMemcpy(devU, hostU, size,  
    cudaMemcpyHostToDevice);  
    example<<<1,4>>>(devU, devV);  
    cudaMemcpy(hostU, devU, size,  
    cudaMemcpyDeviceToHost);  
    cudaFree(devU);  
    return 0;  
}  
  
//Dynamic  
extern __shared__ int tmp[];  
  
__global__ example(float* u) {  
    int i=threadIdx.x;  
     tmp[i]=u[i];  
     u[i]=tmp[i]*tmp[i]+tmp[3-i];  
}  
  
int main() {  
    float hostU[4] = {1, 2, 3, 4};  
    float* devU;  
    size_t size = sizeof(float)*4;  
    cudaMalloc(&devU, size);  
    cudaMemcpy(devU, hostU, size, cudaMemcpyHostToDevice);  
    example<<<1,4,size>>>(devU, devV);  
    cudaMemcpy(hostU, devU, size, cudaMemcpyDeviceToHost);  
    cudaFree(devU);  
    return 0;  
}  

//u(i)= u(i)^2 + u(i-1)
//Static
__global__ example(float* u) {
    int i=threadIdx.x;
    __shared__ int tmp[4];
     tmp[i]=u[i];
     u[i]=tmp[i]*tmp[i]+tmp[3-i];
}

int main() {
    float hostU[4] = {1, 2, 3, 4};
    float* devU;
    size_t size = sizeof(float)*4;
    cudaMalloc(&devU, size);
    cudaMemcpy(devU, hostU, size,
    cudaMemcpyHostToDevice);
    example<<<1,4>>>(devU, devV);
    cudaMemcpy(hostU, devU, size,
    cudaMemcpyDeviceToHost);
    cudaFree(devU);
    return 0;
}

//Dynamic
extern __shared__ int tmp[];

__global__ example(float* u) {
    int i=threadIdx.x;
     tmp[i]=u[i];
     u[i]=tmp[i]*tmp[i]+tmp[3-i];
}

int main() {
    float hostU[4] = {1, 2, 3, 4};
    float* devU;
    size_t size = sizeof(float)*4;
    cudaMalloc(&devU, size);
    cudaMemcpy(devU, hostU, size, cudaMemcpyHostToDevice);
    example<<<1,4,size>>>(devU, devV);
    cudaMemcpy(hostU, devU, size, cudaMemcpyDeviceToHost);
    cudaFree(devU);
    return 0;
}

 4. グローバルメモリ

　　特徴：全スレッドからアクセス可能、キャッシュなし

[cpp］ 表示プレーン コピー

印刷 ?

//Dynamic  
__global__ add4f(float* u, float* v) {  
int i=threadIdx.x;  
 u[i]+=v[i];  
}  
int main() {  
    float hostU[4] = {1, 2, 3, 4};  
    float hostV[4] = {1, 2, 3, 4};  
    float* devU, devV;  
    size_t size = sizeof(float)*4;  
    cudaMalloc(&devU, size);  
    cudaMalloc(&devV, size);  
    cudaMemcpy(devU, hostU, size,  
    cudaMemcpyHostToDevice);  
    cudaMemcpy(devV, hostV, size,  
    cudaMemcpyHostToDevice);  
    add4f<<<<1,4>>>(devU, devV);  
    cudaMemcpy(hostU, devU, size,  
    cudaMemcpyDeviceToHost);  
    cudaFree(devV);  
    cudaFree(devU);  
    return 0;  
}  
  
//static  
__device__ float devU[4];  
__device__ float devV[4];  
  
__global__ addUV() {  
int i=threadIdx.x;  
 devU[i]+=devV[i];  
}  
  
int main() {  
    float hostU[4] = {1, 2, 3, 4};  
    float hostV[4] = {1, 2, 3, 4};  
    size_t size = sizeof(float)*4;  
    cudaMemcpyToSymbol(devU, hostU, size, 0, cudaMemcpyHostToDevice);  
    cudaMemcpyToSymbol(devV, hostV, size, 0, cudaMemcpyHostToDevice);  
     addUV<<<1,4>>>();  
    cudaMemcpyFromSymbol(hostU, devU, size, 0, cudaMemcpyDeviceToHost);  
    return 0;  
}  

//Dynamic
__global__ add4f(float* u, float* v) {
int i=threadIdx.x;
 u[i]+=v[i];
}
int main() {
    float hostU[4] = {1, 2, 3, 4};
    float hostV[4] = {1, 2, 3, 4};
    float* devU, devV;
    size_t size = sizeof(float)*4;
    cudaMalloc(&devU, size);
    cudaMalloc(&devV, size);
    cudaMemcpy(devU, hostU, size,
    cudaMemcpyHostToDevice);
    cudaMemcpy(devV, hostV, size,
    cudaMemcpyHostToDevice);
    add4f<<<<1,4>>>(devU, devV);
    cudaMemcpy(hostU, devU, size,
    cudaMemcpyDeviceToHost);
    cudaFree(devV);
    cudaFree(devU);
    return 0;
}

//static
__device__ float devU[4];
__device__ float devV[4];

__global__ addUV() {
int i=threadIdx.x;
 devU[i]+=devV[i];
}

int main() {
    float hostU[4] = {1, 2, 3, 4};
    float hostV[4] = {1, 2, 3, 4};
    size_t size = sizeof(float)*4;
    cudaMemcpyToSymbol(devU, hostU, size, 0, cudaMemcpyHostToDevice);
    cudaMemcpyToSymbol(devV, hostV, size, 0, cudaMemcpyHostToDevice);
     addUV<<<1,4>>>();
    cudaMemcpyFromSymbol(hostU, devU, size, 0, cudaMemcpyDeviceToHost);
    return 0;
}

5.コンスタント・メモリコンスタント・メモリ

 　　アクセス頻度の高い読み取り専用のパラメータを格納するために使用される

 　　特徴：読み取り専用、キャッシュ、小容量(64KB)

　 　注意：定数メモリを定義する場合、すべての関数の外側で定義する必要があり、ファイル全体に作用します

1 __constant__ int devVar;
2 cudaMemcpyToSymbol(devVar, hostVar, sizeof(int), 0, cudaMemcpyHostToDevice)
3 cudaMemcpyFromSymbol(hostVar, devVar, sizeof(int), 0, cudaMemcpyDeviceToHost)

 6.テクスチャメモリ

     ビデオメモリに1次元、2次元、3次元の配列としてデータを格納した読み出し専用メモリです。汎用コンピューティングにおける画像処理やルックアップの実装に適しており、また、大容量データへのランダムアクセスやアラインドアクセスの高速化にも適している。

     特徴 テクスチャキャッシュを搭載、読み込み専用。

TNE END