Быстрый расчет различий между последовательными индексами - PullRequest
0 голосов
/ 07 февраля 2012

Учитывая, что у меня есть массив

Let Sum be 16
dintptr = { 0 , 2, 8,11,13,15} 

Я хочу вычислить разницу между последовательными индексами, используя графический процессор.Итак, окончательный массив должен быть следующим:

count = { 2, 6,3,2,2,1}

Ниже мое ядро:

//for this function n is 6

__global__ void kernel(int *dintptr, int * count, int n){

   int id = blockDim.x * blockIdx.x + threadIdx.x;
   __shared__ int indexes[256];
   int need = (n % 256 ==0)?0:1;
   int allow = 256 * ( n/256 + need);
   while(id < allow){
     if(id < n ){
       indexes[threadIdx.x] = dintptr[id];

     }
     __syncthreads();
     if(id < n - 1 ){
       if(threadIdx.x % 255 == 0 ){
            count[id] = indexes[threadIdx.x + 1] - indexes[threadIdx.x];
       }else{
            count[id] = dintptr[id+1] - dintptr[id];
       }


    }//end if id<n-1
      __syncthreads();
     id+=(gridDim.x * blockDim.x);
    }//end while
}//end kernel
// For last element explicitly set count[n-1] = SUm - dintptr[n-1]

2 вопроса:

  1. Это ядро ​​быстро.Можете ли вы предложить более быструю реализацию?
  2. Это ядро ​​обрабатывает массивы произвольного размера (я так думаю)

1 Ответ

2 голосов
/ 07 февраля 2012

Я укушу.

__global__ void kernel(int *dintptr, int * count, int n)
{
    for (int id = blockDim.x * blockIdx.x + threadIdx.x; 
         id < n-1; 
         id += gridDim.x * blockDim.x)
        count[id] = dintptr[id+1] - dintptr[i];
}

(Поскольку вы сказали, что «явно» задали значение последнего элемента, а в своем ядре этого не было, я не стал его здесь устанавливать).

Я не вижу большого преимущества использования разделяемой памяти в этом ядре, как вы: кеш L1 на Fermi должен дать вам почти такое же преимущество, так как ваша локальность высока, а повторное использование низкое.

И ваше ядро, и мое, похоже, обрабатывают массивы произвольного размера. Ваш, однако, предполагает, что blockDim.x == 256.

...