Почему версия текстурной памяти нижеприведенной программы медленнее, чем версия глобальной памяти

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;
}

Версия Global Memory:

#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

В дополнение к ответу pQB, в вашей программе нет повторного использования данных - каждый вход читается только один раз и используется только один раз.Индексы памяти являются последовательными для всех потоков и, следовательно, идеально объединены.По этим двум причинам нет необходимости кэшировать память устройства, поэтому глобальный доступ к памяти более эффективен, чем доступ к текстурам.Добавьте к этому дополнительные издержки задержки в кэше текстуры (кэш текстуры предназначен для увеличения пропускной способности, а не уменьшения задержки, в отличие от кэшей данных L1 / L2) и объясните замедление.параллельное сокращение, так что вы можете увидеть пример «сокращения» в CUDA SDK для быстрой реализации.

Answer 2

Хотя известно, что ваше графическое устройство может помочь, для некоторых типов проблем, с вычислительной возможностью 2.x, кэш L1 и L2 работает как кэш текстур.

В этом случае вы не используете кеш текстур, так как вы читаете значение только один раз для потока.С другой стороны, вы используете пространственную локальность в 1D, которую можно скрыть с помощью глобального доступа к памяти.

Я рекомендую вам книгу « CUDA на примере: введение в программирование GPU общего назначения ».Отличная книга для начинающих.С графическими примерами, такими как JuliaSet или очень простой Raycasting (есть также обычные примеры добавления, уменьшения и расстановки точек, если вы предпочитаете thouse :)