Правильно ли проверен пример кода CUDA в gpugems3?

Question

Правильно ли проверен пример кода CUDA в gpugems3?

Я написал кусок кода для вызова ядра в gpugem3

но результаты, которые я получил, - это набор отрицательных чисел вместо префикса сканирования. Мне интересно, если мой вызов ядра неверен или что-то не так с кодом gpugem3?

вот мой код:

#include <stdio.h>
#include <sys/time.h>
#include <cuda.h>



__global__ void kernel(int *g_odata, int  *g_idata, int n, int dim)
{
     extern __shared__ int temp[];// allocated on invocation
    int thid = threadIdx.x;
    int offset = 1;

    temp[2*thid] = g_idata[2*thid]; // load input into shared memory
    temp[2*thid+1] = g_idata[2*thid+1];
    for (int d = n>>1; d > 0; d >>= 1) // build sum in place up the tree
    {
    __syncthreads();
    if (thid < d)
    {
    int ai = offset*(2*thid+1)-1;
    int bi = offset*(2*thid+2)-1;
    temp[bi] += g_idata[ai];
    }
    offset *= 2;
    }
    if (thid == 0) { temp[n - 1] = 0; } // clear the last element
    for (int d = 1; d < n; d *= 2) // traverse down tree & build scan
    {
    offset >>= 1;
    __syncthreads();
    if (thid < d)
    {
    int ai = offset*(2*thid+1)-1;
    int bi = offset*(2*thid+2)-1;
    int t = temp[ai];
    temp[ai] = temp[bi];
    temp[bi] += t;
    }
    }
    __syncthreads();
    g_odata[2*thid] = temp[2*thid]; // write results to device memory
    g_odata[2*thid+1] = temp[2*thid+1];

}


void Initialize(int  *h_in,int num_items)
{

   int j;
   for(j=0;j<num_items;j++)

       h_in[j]=j;
       printf(" input: ");
         printf("\n\n");



}


int main(int argc, char** argv)
{
    int num_items = 512;


    int*  h_in = new int[num_items];


    // Initialize problem 
    Initialize(h_in, num_items);


    int *d_in = NULL;
    cudaMalloc((void**)&d_in, sizeof(int) * num_items);


if(cudaSuccess!=    cudaMemcpy(d_in, h_in, sizeof(int) * num_items, cudaMemcpyHostToDevice)) fprintf(stderr,"could not copy to gpu");

    // Allocate device output array
    int *d_out = NULL;
    cudaMalloc((void**)&d_out, sizeof(int) * (num_items+1));


    kernel<<<1,256,num_items*sizeof(int)>>>(d_out, d_in,num_items, 2);

     int* h_out= new int[num_items+1];
    if( cudaSuccess !=cudaMemcpy(h_out,d_out,sizeof(int)*(num_items+1),cudaMemcpyDeviceToHost))fprintf(stderr,"could not copy back");
    int i;
    printf(" \n");
    for(i=0;i<num_items;i++)
    printf(" ,%d ",h_out[i]);
    // Cleanup
    if (h_in) delete[] h_in;
    if (h_out) delete[] h_out;
    if (d_in) cudaFree(d_in);
    if (d_out) cudaFree(d_out);

    printf("\n\n");

    return 0;
}

0

cuda gpu nvidia prefix-sum

Источник

user3792624 14 июн '15 в 16:52

1 ответ

Решение

Другие вопросы по тегам cuda gpu nvidia prefix-sum

user1695960 14 июн '15 в 22:05 2015-06-14 22:05 · Accepted Answer · 2015-06-14 22:05

Похоже, что вы сделали по крайней мере 1 ошибку при переносе кода из главы GPU Gems 3 в ваше ядро. Эта строка неверна:

temp[bi] += g_idata[ai];

так должно быть:

temp[bi] += temp[ai];

Когда я делаю это одно изменение в коде, который вы сейчас опубликовали, мне кажется, что он выведет для меня правильную сумму префикса (эксклюзивного сканирования). Есть несколько других вещей, которые я бы упомянул:

Даже без этого изменения я получаю некоторые результаты, близкие к правильным. Так что, если вы получаете разные вещи (например, отрицательные числа), у вас могут быть проблемы с настройкой вашего компьютера или установкой CUDA. Я бы предложил использовать более строгую проверку ошибок cuda, чем то, что у вас есть сейчас (хотя проблема с настройкой машины должна была быть указана в одной из ваших проверок).
У рутины, как созданной, будут некоторые ограничения. Он может использоваться только в одном поточном блоке, у него будут конфликты банков при доступе к общей памяти, и он будет ограничен по размеру набора данных тем, что может быть обработано одним потоковым блоком (эта процедура создает два выходных элемента на поток, поэтому ожидается, что размер набора данных будет равен удвоенному числу потоков). Как уже говорилось, динамическое распределение разделяемой памяти должно быть таким же большим, как размер набора данных (т. Е. В два раза больше размера потока в количестве элементов).
Это может быть полезно для обучения, но если вы хотите надежное и быстрое сканирование префиксов, вам рекомендуется использовать подпрограмму из thrust или cub вместо вашего собственного кода, даже если она получена из этой (старой) статьи.

Следующий код похож на ваш, но в нем исправлены вышеуказанные проблемы, и я настроил ядро для использования с различными типами данных:

#include <stdio.h>
#define DSIZE 512
#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)


typedef int mytype;

template <typename T>
__global__ void prescan(T *g_odata, T *g_idata, int n)
{
  extern __shared__ T temp[];  // allocated on invocation
  int thid = threadIdx.x;
  int offset = 1;
  temp[2*thid] = g_idata[2*thid]; // load input into shared memory
  temp[2*thid+1] = g_idata[2*thid+1];
  for (int d = n>>1; d > 0; d >>= 1)                    // build sum in place up the tree
  {
    __syncthreads();
    if (thid < d)
    {
      int ai = offset*(2*thid+1)-1;
      int bi = offset*(2*thid+2)-1;
      temp[bi] += temp[ai];
    }
    offset *= 2;
  }
  if (thid == 0) { temp[n - 1] = 0; } // clear the last element
  for (int d = 1; d < n; d *= 2) // traverse down tree & build scan
    {
      offset >>= 1;
      __syncthreads();
      if (thid < d)
      {
         int ai = offset*(2*thid+1)-1;
         int bi = offset*(2*thid+2)-1;
         T t = temp[ai];
         temp[ai] = temp[bi];
         temp[bi] += t;
      }
    }
  __syncthreads();
  g_odata[2*thid] = temp[2*thid]; // write results to device memory
  g_odata[2*thid+1] = temp[2*thid+1];
}

int main(){

  mytype *h_i, *d_i, *h_o, *d_o;
  int dszp = (DSIZE)*sizeof(mytype);

  h_i = (mytype *)malloc(dszp);
  h_o = (mytype *)malloc(dszp);
  if ((h_i == NULL) || (h_o == NULL)) {printf("malloc fail\n"); return 1;}
  cudaMalloc(&d_i, dszp);
  cudaMalloc(&d_o, dszp);
  cudaCheckErrors("cudaMalloc fail");
  for (int i = 0 ; i < DSIZE; i++){
    h_i[i] = i;
    h_o[i] = 0;}
  cudaMemset(d_o, 0, dszp);
  cudaCheckErrors("cudaMemset fail");
  cudaMemcpy(d_i, h_i, dszp, cudaMemcpyHostToDevice);
  cudaCheckErrors("cudaMemcpy 1 fail");
  prescan<<<1,DSIZE/2, dszp>>>(d_o, d_i, DSIZE);
  cudaDeviceSynchronize();
  cudaCheckErrors("kernel fail");
  cudaMemcpy(h_o, d_o, dszp, cudaMemcpyDeviceToHost);
  cudaCheckErrors("cudaMemcpy 2 fail");
  mytype psum = 0;
  for (int i =1; i < DSIZE; i++){
    psum += h_i[i-1];
    if (psum != h_o[i]) {printf("mismatch at %d, was: %d, should be: %d\n", i, h_o[i], psum); return 1;}
    }
  return 0;
}