なぜテクスチャメモリのバージョンがグローバルメモリのバージョンよりも遅い

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キャッシュは、テクスチャキャッシュとして良好に機能します。

この場合、スレッドごとに値を1度だけ読み取るので、テクスチャキャッシュを利用していません。もう1つのチャンドでは、1Dの空間的局所性を利用して、グローバルメモリ統合アクセスで非表示にすることができます。

私はあなたに書籍「CUDA by Example：General Purpose GPU Programming」をお勧めします。初心者のための素晴らしい本。 JuliaSetや非常に基本的なRaycastingのようなグラフィックスの例もあります（あなたが好きな人は、共通のadd、reduce、dotの例があります:)。

このヘルプが必要です。

Answer 2

さらに、pQBの答えによれば、プログラムのデータ再利用はありません。各入力は一度だけ読み込まれ、一度だけ使用されます。メモリインデックスはスレッド間で連続しているため、完全に結合されています。この2つの理由から、デバイスメモリのキャッシュは必要ありません。したがって、グローバルメモリアクセスはテクスチャアクセスよりも効率的です。テクスチャキャッシュのこの追加レイテンシオーバヘッドに加えて（テクスチャキャッシュは、L1/L2データキャッシュとは異なり、スループットを向上させ、レイテンシを減少させないように設計されています）、スローダウンについて説明します。

ところで、あなたがやっていることは並行処理の削減です。そのため、CUDA SDKの「削減」の例を高速に実装したい場合があります。