Индексирование отрицательного массива на основе 1d трафарета на основе общей памяти Реализация CUDA

Question

В настоящее время я работаю с программированием CUDA, и я пытаюсь изучить слайды из семинара, который я нашел в Интернете, который можно найти here. Проблема, которая возникает у меня на слайде 48. Здесь можно найти следующий код:

__global__ void stencil_1d(int *in, int *out) { 

    __shared__ int temp[BLOCK_SIZE + 2 * RADIUS]; 

    int gindex = threadIdx.x + blockIdx.x * blockDim.x; 
    int lindex = threadIdx.x + RADIUS; 

    // Read input elements into shared memory 
    temp[lindex] = in[gindex]; 
    if (threadIdx.x < RADIUS) { 
     temp[lindex - RADIUS] = in[gindex - RADIUS]; 
     temp[lindex + BLOCK_SIZE] = in[gindex + BLOCK_SIZE]; 
    } 

.... 

Чтобы добавить немного контекста. У нас есть массив под названием in, который, как говорят, N. Затем мы имеем еще один массив out, который имеет длину N+(2*RADIUS), где RADIUS имеет значение 3 для данного конкретного примера. Идея состоит в том, чтобы скопировать массив in в массив out, но поместить массив in в положение 3 с начала массива out i.e out = [RADIUS][in][RADIUS], см. Слайд для графического представления.

Путаница происходит в на следующей строке:

temp[lindex - RADIUS] = in[gindex - RADIUS]; 

Если gindex является 0 то мы имеем in[-3]. Как мы можем читать из отрицательного индекса в массиве? Любая помощь будет действительно оценена.

Answer 1

Ответ на вопрос pQB верен. Вы должны смещать указатель входного массива на RADIUS.

Чтобы показать это, я предоставляю ниже полный рабочий пример. Надеюсь, что это будет полезно для других пользователей.

(я бы сказал, что после загрузки разделяемой памяти вам понадобится __syncthreads(). Я добавил его в приведенном ниже примере).

#include <thrust/device_vector.h> 

#define RADIUS  3 
#define BLOCKSIZE 32 

/*******************/ 
/* iDivUp FUNCTION */ 
/*******************/ 
int iDivUp(int a, int b){ return ((a % b) != 0) ? (a/b + 1) : (a/b); } 

/********************/ 
/* CUDA ERROR CHECK */ 
/********************/ 
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } 
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true) 
{ 
    if (code != cudaSuccess) 
    { 
     fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); 
     if (abort) exit(code); 
    } 
} 

/**********/ 
/* KERNEL */ 
/**********/ 
__global__ void moving_average(unsigned int *in, unsigned int *out, unsigned int N) { 

    __shared__ unsigned int temp[BLOCKSIZE + 2 * RADIUS]; 

    unsigned int gindexx = threadIdx.x + blockIdx.x * blockDim.x; 

    unsigned int lindexx = threadIdx.x + RADIUS; 

    // --- Read input elements into shared memory 
    temp[lindexx] = (gindexx < N)? in[gindexx] : 0; 
    if (threadIdx.x < RADIUS) { 
     temp[threadIdx.x] = (((gindexx - RADIUS) >= 0)&&(gindexx <= N)) ? in[gindexx - RADIUS] : 0; 
     temp[threadIdx.x + (RADIUS + BLOCKSIZE)] = ((gindexx + BLOCKSIZE) < N)? in[gindexx + BLOCKSIZE] : 0; 
    } 

    __syncthreads(); 

    // --- Apply the stencil 
    unsigned int result = 0; 
    for (int offset = -RADIUS ; offset <= RADIUS ; offset++) { 
     result += temp[lindexx + offset]; 
    } 

    // --- Store the result 
    out[gindexx] = result; 
} 

/********/ 
/* MAIN */ 
/********/ 
int main() { 

    const unsigned int N  = 55 + 2 * RADIUS; 

    const unsigned int constant = 4; 

    thrust::device_vector<unsigned int> d_in(N, constant); 
    thrust::device_vector<unsigned int> d_out(N); 

    moving_average<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(thrust::raw_pointer_cast(d_in.data()), thrust::raw_pointer_cast(d_out.data()), N); 
    gpuErrchk(cudaPeekAtLastError()); 
    gpuErrchk(cudaDeviceSynchronize()); 

    thrust::host_vector<unsigned int> h_out = d_out; 

    for (int i=0; i<N; i++) 
     printf("Element i = %i; h_out = %i\n", i, h_out[i]); 

    return 0; 

} 

Answer 2

Вы считаете, что массив in указывает на первую позицию памяти, выделенную для этого массива. Однако, если вы видите слайд 47, массив in имеет гало (оранжевые ящики) из трех элементов до и после данных (показаны как зеленые кубики).

Мое предположение (я не делал мастерскую), что входной массив сначала инициализируется ореолом, а затем указатель перемещается в ядре. Что-то вроде:

stencil_1d<<<dimGrid, dimBlock>>>(in + RADIUS, out); 

Таким образом, в ядре, это безопасно in[-3], потому что указатель не на начало массива.

Answer 3

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

В C (не только в CUDA, но в C вообще), когда вы получаете доступ к «массиву «с помощью скобок [], вы на самом деле делаете указатель арифметики.

Для примера рассмотрим указатель, как это:

int* data= ... // Points to some memory 

Когда вы затем написать заявление, как

data[3] = 42; 

вы только доступ к ячейке памяти, которая «три записи за оригинальный data указатель". Таким образом, вы можете также написали

int* data= ... // Points to some memory 
int* dataWithOffset = data+3; 
dataWithOffset[0] = 42; // This will write into data[3] 

и, следовательно,

dataWithOffset[-3] = 123; // This will write into data[0] 

---

_{На самом деле, вы можете сказать, что data[i] такое же, как *(data+i), который является таким же, как *(i+data), что в поворот такой же, как и i[data], но вы не должны использовать его в реальных программах ...)}