CUDA generic dot product

Question

Я новичок как в C, так и в CUDA, и я писал функцию продукта dot, однако это не дает мне правильных результатов. Будет ли какая-то добродушная душа помочь мне взглянуть?

У меня также есть два вопроса,

1. Почему точка() не работает правильно, и

2. В строке 57, почему это произведение [threadIdx.x] вместо продукта [индекс]? Могу ли я не писать

   продукт [индекс] = a [указатель] * b [индекс]; ... if (index == 0) {...} и суммировать каждый элемент с нулевой нитью таким образом?

Большое спасибо.

deviceQuery:

Device 0: "GeForce GTX 570" 
    CUDA Driver Version/Runtime Version   6.0/5.5 
    CUDA Capability Major/Minor version number: 2.0 

Makefile: NVCC -arch = sm_20 cuda_test.cu -o cuda_test

в cuda_test.cu:

#include <stdio.h> // printf, scanf, NULL etc. 
#include <stdlib.h> // malloc, free, rand etc. 

#define N (3) //Number of threads we are using (also, length of array declared in main) 

#define THREADS_PER_BLOCK (1) //Threads per block we are using 

#define N_BLOCKS (N/THREADS_PER_BLOCK) 

/* Function to generate a random integer between 1-10 */ 
void random_ints (int *a, int n) 
{ 
    int i; 
    srand(time(NULL)); //Seed rand() with current time 
    for(i=0; i<n; i++) 
    { 
     a[i] = rand()%10 + 1; 
    } 
    return; 
} 

/* Kernel that adds two integers a & b, stores result in c */ 
__global__ void add(int *a, int *b, int *c) { 
//global indicates function that runs on 
//device (GPU) and is called from host (CPU) code 

    int index = threadIdx.x + blockIdx.x * blockDim.x; 

    //threadIdx.x : thread index 
    //blockIdx.x : block index 
    //blockDim.x : threads per block 
    //hence index is a thread counter across all blocks 
    c[index] = a[index] + b[index]; 

//note that pointers are used for variables 
//add() runs on device, so they must point to device memory 
//need to allocate memory on GPU 
} 

/* Kernel for dot product */ 
__global__ void dot(int *a, int *b, int *c) 
{ 
    __shared__ int product[THREADS_PER_BLOCK]; //All threads in a block must be able 
               //to access this array 

    int index = threadIdx.x + blockIdx.x * blockDim.x; //index 

    product[threadIdx.x] = a[index] * b[index]; //result of elementwise 
               //multiplication goes into product 

    //Make sure every thread has finished 
    __syncthreads(); 

    //Sum the elements serially to obtain dot product 
    if(0 == threadIdx.x) //Pick one thread to sum, otherwise all will execute 
    { 
     int sum = 0; 
     for(int j=0; j < THREADS_PER_BLOCK; j++) sum += product[j]; 
     //Done! 
     atomicAdd(c,sum); 
    } 
} 

int main(void) 
{ 

    int *a, *b, *c, *dotProduct; //host copies of a,b,c etc 
    int *d_a, *d_b, *d_c, *d_dotProduct; //device copies of a,b,c etc 

    int size = N * sizeof(int); //size of memory that needs to be allocated 

    int i=0; //iterator 

    //Allocate space for device copies of a,b,c 
    cudaMalloc((void **)&d_a, size); 
    cudaMalloc((void **)&d_b, size); 
    cudaMalloc((void **)&d_c, size); 

    //Setup input values 
    a = (int *)malloc(size); random_ints(a,N); 
    b = (int *)malloc(size); random_ints(b,N); 
    c = (int *)malloc(size); 

    //Copy inputs to device 
    cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice); 
    cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice); 

    //Launch add() kernel on GPU 
    add<<<N_BLOCKS,THREADS_PER_BLOCK>>>(d_a, d_b, d_c); 
    // triple angle brackets mark call from host to device 
    // this is also known as a kernel launch 
    // N/THREADS_PER_BLOCK = NO. OF BLOCKS 

    //Copy result back to host 
    cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost); 

    //Output results 
    printf("a = {"); 
    for (i=0; i<N; i++) printf(" %d",a[i]); 
    printf(" }\n"); 

    printf("b = {"); 
    for (i=0; i<N; i++) printf(" %d",b[i]); 
    printf(" }\n"); 

    printf("c = {"); 
    for (i=0; i<N; i++) printf(" %d",c[i]); 
    printf(" }\n"); 

    //Calculate dot product of a & b 
    dotProduct = (int *)malloc(sizeof(int)); //Allocate host memory to dotProduct 
    *dotProduct = 0; //initialise to zero 
    cudaMalloc((void **)&d_dotProduct, sizeof(int)); //Allocate device memory to d_dotProduct 
    dot<<<N_BLOCKS,THREADS_PER_BLOCK>>>(d_a, d_b, d_dotProduct); //Perform calculation 
    cudaMemcpy(dotProduct, d_dotProduct, sizeof(int), cudaMemcpyDeviceToHost); //Copy result into dotProduct 
    printf("\ndot(a,b) = %d\n", *dotProduct); //Output result 

    //Cleanup 
    free(a); free(b); free(c); free(dotProduct); 
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); cudaFree(d_dotProduct); 

    return 0; 
} //End of main 

Answer 1

Как talonmies говорит, пожалуйста, сделать так, кто-то может запустить ваш код. Внедрение номеров строк бесполезно.

Лучшая догадка без другой информации заключается в том, что вы не инициализировали d_dotProduct до нуля. Вы можете сделать это с помощью cudaMemset() - если бы вы хотели получить другое начальное значение, вы могли бы cudaMemcpy() получить начальное значение от хоста или запустить отдельное ядро ​​для инициализации, но в этом случае достаточно cudaMemset() (что эквивалентно на хосте).

Также может быть N_BLOCKS*THREADS_PER_BLOCK не равно size.

Что касается вашего второго вопроса, то product представляет собой массив блоков размером THREADS_PER_BLOCK, если вы обращаетесь к нему с помощью product[index], вы окажетесь вне пределов.

Answer 2

проблема решен! необходимо было установить «* c = 0» перед суммированием отдельных элементов массива «продукт».

/* Kernel for dot product */ 
__global__ void dot(int *a, int *b, int *c) 
{ 
    __shared__ int product[THREADS_PER_BLOCK]; //All threads in a block must be able 
               //to access this array 

    int index = threadIdx.x + blockIdx.x * blockDim.x; //index 

    product[threadIdx.x] = a[index] * b[index]; //result of elementwise 
               //multiplication goes into product 

    if(index==0) *c = 0; //Ask one thread to set c to zero. 

    //Make sure every thread has finished 
    __syncthreads();  

    //Sum the elements serially to obtain dot product 
    if(0 == threadIdx.x) //Every block to do c += sum 
    { 
     int sum = 0; 
     for(int j=0; j < THREADS_PER_BLOCK; j++) sum += product[j]; 
     //Done! 
     atomicAdd(c,sum); 
    } 
}