Как мне скопировать входной массив cuda в общий массив?

#cuda #shared-memory

#cuda #разделяемая память

Вопрос:

Я пытаюсь скопировать входной массив cuda в массив общей памяти. Первые n значений отлично копируются в общий массив, но после этого происходят некоторые довольно странные закономерности. Кто-нибудь может найти, что я делаю неправильно / что здесь происходит?

Я пытался использовать s_x[tid] = x [tid], s_x [row] = x [row], а также две другие комбинации. Я знаю, что они логически не имеют смысла, но нужно попробовать все.

Вот функция ядра, которую я написал:

     __global__ void mv_cuda_shared(float* y, float* A, float* x, int n)
    {
        extern __shared__ float s_x[];
        int row = blockIdx.x * blockDim.x   threadIdx.x;
        int tid = threadIdx.x;
        s_x[row] = x[row];
        __syncthreads();
        if(tid == 0 amp;amp; row == 0)
                for(int i = 0; i < n; i  ){printf("s_x[%d] = .6f x = .6fn", i, s_x[i], x[i]);}
        float temp = 0;
        if(row < n)
        {
            for(int k = 0; k < n; k  )
            {
                temp  = (A[row*n   k] * s_x[k]);
            }
        }
        y[row] = temp;
    }
  

и здесь я вызываю его в основном методе (опуская все cuda mallocs и копируя):

     mv_cuda_shared<<<(N/BLOCK_SIZE), BLOCK_SIZE, N*sizeof(float)>>>(d_y, d_A, d_x, N);
  

Поэтому я ожидаю, что s_x будет соответствовать x. Но вместо этого первые N (n = N) элементов копируются правильно, но затем остальные, похоже, следуют шаблону. Например, 4 записи после s_x [N-1], скажем, равны 1. Тогда 4 после этого похожи на .24566…. (Я проверю, и число даже не находится в исходном массиве). Итак, кто-нибудь может увидеть, что я сделал не так?

Комментарии:

1. s_x[row] определенно неправильно. N*sizeof(float) тоже неправильно, так и должно быть BLOCK_SIZE*sizeof(float) . Вы запускаете только N потоки в целом. Почему вы ожидаете, что будет скопировано больше, чем N значения (например, почему вы ожидаете, что значения после s_x[N-1] будут иметь какое-либо значение)?

Ответ №1:

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

     extern __shared__ float s_x[];
    int tid = threadIdx.x;
    for(int i=tid; i<n; i =blockDim.x) {
       s_x[i] = x[i];
    }
    __syncthreads();
  

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