定数メモリとグローバルメモリを使用するプログラムの違い

Question

私には2つのプログラムがあります。唯一の違いは、定数メモリを使用して入力を格納し、もう一方はグローバルメモリを使用することです。グローバルメモリが定数メモリよりも高速である理由を知りたいのですが？それらは両方とも、内積のbtw 2行列を計算します。

#include<cuda_runtime.h> 
#include<cuda.h> 
#include<stdio.h> 
#include<stdlib.h> 
#define intMin(a,b) ((a<b)?a:b) 
//Threads per block 
#define TPB 128 
//blocks per grid 
#define BPG intMin(128, ((n+TPB-1)/TPB)) 

const int n = 4; 
__constant__ float deva[n],devb[n]; 
__global__ void addVal(float *c){ 
    int tid = blockIdx.x * blockDim.x + threadIdx.x; 

    //Using shared memory to temporary store results 
    __shared__ float cache[TPB]; 
    float temp = 0; 
    while(tid < n){ 
     temp += deva[tid] * devb[tid]; 
     tid += gridDim.x * blockDim.x; 


    } 
    cache[threadIdx.x] = temp; 
    __syncthreads(); 
    int i = blockDim.x/2; 
    while(i !=0){ 
     if(threadIdx.x < i){ 
      cache[threadIdx.x] = cache[threadIdx.x] +cache[threadIdx.x + i] ; 

     } 
    __syncthreads(); 
    i = i/2; 

    } 
    if(threadIdx.x == 1){ 
     c[blockIdx.x ] = cache[0]; 
    } 



} 



int main(){ 

float a[n] , b[n] , c[BPG]; 
//float *deva, *devb, *devc; 
float *devc; 
int i; 
//Filling with random values to test 
for(i =0; i< n; i++){ 
    a[i] = i; 
    b[i] = i*2; 
} 

//cudaMalloc((void**)&deva, n * sizeof(float)); 
//cudaMalloc((void**)&devb, n * sizeof(float)); 

cudaMalloc((void**)&devc, BPG * sizeof(float)); 
//cudaMemcpy(deva, a, n *sizeof(float), cudaMemcpyHostToDevice); 
//cudaMemcpy(devb, b, n *sizeof(float), cudaMemcpyHostToDevice); 
cudaMemcpyToSymbol(deva, a, n * sizeof(float)); 
cudaMemcpyToSymbol(devb, b, n * sizeof(float)); 
cudaEvent_t start, stop; 
cudaEventCreate(&start); 
cudaEventCreate(&stop); 
cudaEventRecord(start, 0); 

//Call function to do dot product 
addVal<<<BPG, TPB>>>(devc); 
cudaEventRecord(stop, 0); 
cudaEventSynchronize(stop); 
float time; 
cudaEventElapsedTime(&time,start, stop); 
printf("The elapsed time is: %f\n", time); 

//copy result back 
cudaMemcpy(c, devc, BPG * sizeof(float), cudaMemcpyDeviceToHost); 
float sum =0 ; 
for (i = 0 ; i< BPG; i++){ 
    sum+=c[i]; 

} 
//display answer 
printf("%f\n",sum); 


getchar(); 

return 0; 
} 

以下はグローバルメモリのバージョンです。

#include<cuda_runtime.h> 
#include<cuda.h> 
#include<stdio.h> 
#include<stdlib.h> 
#define intMin(a,b) ((a<b)?a:b) 
//Threads per block 
#define TPB 128 
//blocks per grid 
#define BPG intMin(128, ((n+TPB-1)/TPB)) 

const int n = 4; 

__global__ void addVal(float *a, float *b, float *c){ 
    int tid = blockIdx.x * blockDim.x + threadIdx.x; 

    //Using shared memory to temporary store results 
    __shared__ float cache[TPB]; 
    float temp = 0; 
    while(tid < n){ 
     temp += a[tid] * b[tid]; 
     tid += gridDim.x * blockDim.x; 


    } 
    cache[threadIdx.x] = temp; 
    __syncthreads(); 
    int i = blockDim.x/2; 
    while(i !=0){ 
     if(threadIdx.x < i){ 
      cache[threadIdx.x] = cache[threadIdx.x] +cache[threadIdx.x + i] ; 

     } 
    __syncthreads(); 
    i = i/2; 

    } 
    if(threadIdx.x == 1){ 
     c[blockIdx.x ] = cache[0]; 
    } 



} 

int main(){ 

float a[n] , b[n] , c[BPG]; 
float *deva, *devb, *devc; 
int i; 
//Filling with random values to test 
for(i =0; i< n; i++){ 
    a[i] = i; 
    b[i] = i*2; 
} 
printf("Not using constant memory\n"); 
cudaMalloc((void**)&deva, n * sizeof(float)); 
cudaMalloc((void**)&devb, n * sizeof(float)); 
cudaMalloc((void**)&devc, BPG * sizeof(float)); 
cudaMemcpy(deva, a, n *sizeof(float), cudaMemcpyHostToDevice); 
cudaMemcpy(devb, b, n *sizeof(float), cudaMemcpyHostToDevice); 

cudaEvent_t start, stop; 
cudaEventCreate(&start); 
cudaEventCreate(&stop); 
cudaEventRecord(start, 0); 

//Call function to do dot product 
addVal<<<BPG, TPB>>>(deva, devb, devc); 
cudaEventRecord(stop, 0); 
cudaEventSynchronize(stop); 
float time; 
cudaEventElapsedTime(&time,start, stop); 
printf("The elapsed time is: %f\n", time); 


//copy result back 
cudaMemcpy(c, devc, BPG * sizeof(float), cudaMemcpyDeviceToHost); 
float sum =0 ; 
for (i = 0 ; i< BPG; i++){ 
    sum+=c[i]; 

} 
//display answer 
printf("%f\n",sum); 


getchar(); 

return 0; 
} 

Answer 1

あなたは定数メモリを利用していません。

• 定数メモリからの1回の読み取りは、ハーフワープにブロードキャストできます（独自のtidからすべてのスレッドがロードされるため、あなたのケースではありません）。
• 定数メモリはキャッシュされます（定数メモリ配列の各位置から一度だけ読み込むので、あなたのケースでは使用されません）。

ハーフワープ内の各スレッドが異なるデータに対して1回の読み取りを実行するため、16回の異なる読み取りがシリアル化され、要求を配置する時間の16倍になります。

グローバルメモリから読み取っている場合、要求は同時に行われます。が合体してになります。そのため、グローバルメモリの例が定数メモリよりも優れています。

もちろん、この結論は、L1およびL2キャッシュを備えた計算機能2.xのデバイスによって異なる可能性があります。

よろしくお願いいたします。