爲什麼下面的程序比全局內存版

Question

我很困惑，爲什麼我的紋理版本比我的全局內存的版本更慢，因爲質地版本應該充分利用空間局部性慢紋理內存版本。我正試圖在下面的情況下計算點積。因此，如果一個線程訪問索引i，則其鄰居應訪問i + 1。因此，我們看到空間局部性。

下面是紋理內存版本：

#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)) 

texture<float> arr1; 
texture<float> arr2; 


const int n = 4; 

__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 += tex1Dfetch(arr1,tid) * tex1Dfetch(arr2,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); 
cudaBindTexture(NULL,arr1, deva,sizeof(float) * n); // note: deva shd be in gpu 
cudaBindTexture(NULL,arr2, devb,sizeof(float) * n); // note: deva shd be in gpu 
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); 
cudaUnbindTexture(arr1); 
cudaUnbindTexture(arr2); 
cudaFree(devc); 

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

雖然知道自己的圖形設備可能會有幫助，對於某些類型的問題，具有計算能力的2.X L1和L2緩存與紋理緩存一樣好。

在這種情況下，你是不是利用紋理緩存，因爲你只有一次每個線程值讀取。另一方面，您正在利用1D中的空間局部性，可以通過全局內存聯合訪問來隱藏。

我推薦你這本書'CUDA by Example：通用GPU編程介紹'。適合初學者的好書。隨着顯卡的例子像JuliaSet或一個非常基本的光線投射（也有共同的添加，減少和點積的例子，如果你喜歡thouse :)。

希望得到這個幫助。

Answer 2

繼PQB的答案，還有在你的程序中沒有數據重用 - 每個輸入只讀一次，並且只能使用一次。內存索引在各個線程之間是順序的，因此完全合併。由於這兩個原因，不需要任何設備內存緩存，所以全局內存訪問比紋理訪問更有效。添加到在紋理高速緩存此額外的等待時間開銷（紋理高速緩存被設計爲提高吞吐量，不會降低延遲，不同於L1/L2數據高速緩存），並且減速進行說明。

順便說一句，你在做什麼是一個平行的減少，所以你可能希望看到一個快速實施的CUDA SDK中的「減量化」的例子。