2011-07-09 80 views
4

我有兩個程序。唯一的區別是使用常量內存來存儲輸入,而另一個使用全局內存。我想知道爲什麼全局內存比常量內存更快?他們都計算dot產品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; 
} 
+0

對我來說,你看到的兩個版本完全相同。所以你可能想要檢查你是否有兩個版本。如果我錯誤地認爲兩個版本是相同的,那麼如果您能夠突出顯示差異在哪裏,可以更容易地找到它們,這將非常有幫助。此外,不同世代的cuda設備在性能特徵方面差異很大,因此,如果您告訴我們,您可能會遇到這種情況(「可能」),因爲我不記得cuda常數的細節記憶,我不知道) – Grizzly

+0

謝謝fr指出,花花公子。我編輯了代碼 – Programmer

回答

9

你沒有獲得固定內存的優勢。

  • 從常量內存中的單個讀取可以廣播到半扭曲(不是你的情況,因爲每個線程從它自己的tid中加載)。
  • 常量內存被緩存(在您的情況下不使用,因爲您只能從常量內存數組中的每個位置讀取一次)。

由於半扭曲中的每個線程都會對不同的數據進行單獨讀取,因此16個不同的讀取將被序列化,這需要16倍的時間來發出請求。

如果他們正在從全局內存中讀取,則請求將在同一時間完成,合併。這就是爲什麼你的全局內存例子比常量內存更好。

當然,這個結論會隨着具有L1和L2緩存的計算能力2.x的設備而變化。

問候!

+2

對於計算能力2.0(sm_20)GPU及更高版本,用「warp」替換上面的「half-warp」,用「32」替換「16」。 – harrism

+0

謝謝你的評論:)。問候! – pQB

相關問題