тепло матричного уравнения в CUDA - нелегальный ошибка адреса

Question

После this question со ссылкой на shared memory example в официальном руководстве, я пытаюсь построить матрицу уравнения теплопроводности, которое так же, как в этом плохо нарисованном изображении, которое я сделал

Вот что я сделал до сих пор, минимальный пример

#define N 32 
#define BLOCK_SIZE 16 
#define NUM_BLOCKS ((N + BLOCK_SIZE - 1)/ BLOCK_SIZE) 

__global__ void heat_matrix(int* A) 
{ 
    const unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x; 
    __shared__ int temp_sm_A[N*N]; 
    int* temp_A = &temp_sm_A[0]; memset(temp_A, 0, N*N*sizeof(int)); 

    if (tid < N) //(*) 
    { 
     #pragma unroll 
     for (unsigned int m = 0; m < NUM_BLOCKS; ++m) 
     {   
      #pragma unroll 
      for (unsigned int e = 0; e < BLOCK_SIZE ; ++e) 
      { 
       if ((tid == 0 && e == 0) || (tid == (N-1) && e == (BLOCK_SIZE-1))) 
       { 
        temp_A[tid + (e + BLOCK_SIZE * m) * N] = -2; 
        temp_A[tid + (e + BLOCK_SIZE * m) * N + (tid==0 ? 1 : -1)] = 1; 
       } 
       if (tid == e) 
       { 
        temp_A[tid + (e + BLOCK_SIZE * m) * N - 1] = 1; 
        //printf("temp_A[%d] = 1;\n", (tid + (e + BLOCK_SIZE * m) * N -1)); 
        temp_A[tid + (e + BLOCK_SIZE * m) * N] = -2; 
        //printf("temp_A[%d] = -2;\n", (tid + (e + BLOCK_SIZE * m) * N)); 
        temp_A[tid + (e + BLOCK_SIZE * m) * N + 1] = 1; 
        //printf("temp_A[%d] = 1;\n", (tid + (e + BLOCK_SIZE * m) * N +1)); 
       } 
      } 
     } 
     __syncthreads(); //(**) 
     memcpy(A, temp_A, N*N*sizeof(int)); 
    } 
} 
int main(){ 
    int* h_A = (int*)malloc(N*N*sizeof(int)); memset(h_A, 0, N*N*sizeof(int)); 
    int* d_A; 
    checkCudaErrors(cudaMalloc((void**)&d_A, N*N*sizeof(int))); 
    checkCudaErrors(cudaMemcpy(d_A, h_A, N*N*sizeof(int), cudaMemcpyHostToDevice)); 
    dim3 dim_grid((N/2 + BLOCK_SIZE -1)/ BLOCK_SIZE); 
    dim3 dim_block(BLOCK_SIZE); 

    heat_matrix <<< dim_grid, dim_block >>> (d_A); 
    checkCudaErrors(cudaMemcpy(h_A, d_A, N*N*sizeof(int), cudaMemcpyDeviceToHost)); 
... 
} 

код, расположенный в соответствии с большой N (больше 32). Я воспользовался блочным разделением. При выполнении nvcc дает

CUDA error at matrix.cu:102 code=77(cudaErrorIllegalAddress) "cudaMemcpy(h_A, d_A, N*N*sizeof(int), cudaMemcpyDeviceToHost)" 

cuda-memcheck И обеспечивает только одна ошибка (на самом деле есть еще один, но это происходит от cudasuccess=checkCudaErrors(cudaDeviceReset()); ...)

========= CUDA-MEMCHECK 
========= Invalid __shared__ write of size 4 
=========  at 0x00000cd0 in heat_matrix(int*) 
=========  by thread (0,0,0) in block (0,0,0) 
=========  Address 0xfffffffc is out of bounds 
... 

Я не могу видеть, где я сделал неправильно в коде. Как может поток 0 в первом блоке спровоцировать незаконный доступ? Есть даже конкретный случай if, чтобы справиться с ним, и не сообщается строка кода, в которой произошла ошибка.

Кроме того, есть ли более эффективный способ для моего кода, чем иметь дело со всеми этими if s? Конечно, есть, но я не мог найти лучшего параллельного выражения, чтобы разделить случаи на второй for.

---

На стороне записки, мне (*) кажется ненужным; вместо этого (**) необходим, если я хочу следовать другим вызовам функций GPU. Я прав?

Answer 1

1. проверить эту строку:

     temp_A[tid + (e + BLOCK_SIZE * m) * N - 1] = 1; 
   

   для потока с tid равна нулю в течение первой итерации, tid + (e + BLOCK_SIZE * m) * N - 1 имеет значение индекса -1. Это то, о чем жалуется выпуск cuda-memcheck (с адресом, который был обернут из-за недостаточного потока).

2. Похожая вне границ доступа будет происходить позже для линии

     temp_A[tid + (e + BLOCK_SIZE * m) * N + 1] = 1; 
   

   когда tid, e и m все предположить их максимальное значение.

3. У вас есть несколько потоков, пишущих в том же месте памяти. Каждый поток должен писать ровно один элемент массива на каждую итерацию внутреннего цикла. Нет необходимости выписывать соседние элементы, потому что они уже покрыты своими потоками.

4. У вас есть условие гонки между инициализацией memset() и магазинами внутри основных петель. Поместите syncthreads() после memset().

5. Призывы к memset() и memcpy() приведет к каждому потоку делать полный набор/копировать, делать операции N раз вместо того, чтобы только один раз.
   Общим способом обработки этого является выписать операцию явно, разделив работу между потоками блока.
   Однако ...

6. нет никакой пользы от создания матрицы в разделяемой памяти, а затем копирования ее в глобальную память позже. Запись в A в глобальной памяти устраняет необходимость в memset(), memcpy() и syncthreads() в целом.

7. Использование размера блока всего лишь из 16 потоков оставляет половину ресурсов неиспользованными, так как блоки потоков распределены в единицах из 32 потоков (деформация).

Вы можете перечитать раздел о Thread Hierarchy в Руководстве по программированию CUDA C.

Answer 2

В вашем ядре temp_A является локальным указателем на начало вашего массива разделяемой памяти. С учетом:

N = 32;

BLOCK_SIZE = 16;

m (0,1);

е (0, BLOCK_SIZE)

Обращается как temp_A[tid + (e + BLOCK_SIZE * m) * N] может легко выйти из границ 1024-элементов длиной массива.