Ошибка общей памяти pycuda "pycuda._driver.LogicError: ошибка cuLaunchKernel: недопустимое значение"

Question

У меня странная проблема, происхождение которой я не могу определить:

У меня есть работающее ядро ​​для некоторого специального умножения матриц-векторов, которое я хочу ускорить. В основном большая матрица (10^6 умноженная на 10^6) построена из нескольких маленьких матриц. Поэтому я хочу поместить эти данные в общую память. Однако, когда я пытаюсь добавить общую память, я получаю только ошибку:

pycuda._driver.LogicError: ошибка cuLaunchKernel: недопустимое значение

Итак, мое рабочее ядро:

#define FIELD_SIZE {field}
#define BLOCK_SIZE {block}

__global__ void MatrixMulKernel(double *gpu_matrix, double *gpu_b, double *gpu_y)
{
    int tx = ... + threadIdx.x;

    if(tx < FIELD_SIZE*FIELD_SIZE*BLOCK_SIZE) 
    { ... multiplication ... }
}

И если я пытаюсь добавить часть общей памяти, это выглядит так

#define FIELD_SIZE {field}
#define BLOCK_SIZE {block}

__global__ void MatrixMulKernel(double *gpu_matrix_ptr, double *gpu_b, double *gpu_y)
{
    __shared__ double gpu_matrix[BLOCK_SIZE*BLOCK_SIZE*13];

    int tx = ... + threadIdx.x;
    if(tx < BLOCK_SIZE*BLOCK_SIZE*13) {  gpu_matrix[tx] = gpu_matrix_ptr[tx];  }
    __syncthreads();

    if(tx < FIELD_SIZE*FIELD_SIZE*BLOCK_SIZE) 
    { ... multiplication ... }
}

Это единственная часть, которую я изменил, так что в основном это должен быть оператор gpu_matrix[tx] = gpu_matrix_ptr[tx], не так ли? Но я не понимаю, как это должно быть. Я в основном пытался скопировать пример умножения мозаичных матриц из примеров Pycuda. http://wiki.tiker.net/PyCuda/Examples/MatrixmulTiled

Вызов:

self.kernel.prepare([np.intp, np.intp, np.intp])
self.kernel.prepared_call(grid_shape,
              block_shape,
              self.matrix_gpu.gpudata,
              b_gpu.gpudata,
              y_gpu.gpudata)

где matrix_gpu, b_gpu и y_gpu являются экземплярами pycuda.gpuarray.

Надеюсь, что вы можете прояснить некоторые мои замешательства...

Answer 1

Согласно вашему описанию, общий доступ к памяти слишком велик.

__shared__ double gpu_matrix[BLOCK_SIZE*BLOCK_SIZE*13];

shared mem является одним из аппаратных ресурсов cuda gpu. общий размер около 48Кбайт, который вы не можете увеличить.

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

$CUDA_ROOT/tools/CUDA_Occupancy_Calculator.xls

С другой стороны, размер разделяемой памяти, требуемый ядрами mat-vec-mul-like, должен быть в состоянии уменьшиться с O (BLOCK_SIZE^2) тоже(BLOCK_SIZE). Возможно, вы захотите прочитать код некоторых успешных ядер mat-vec-mul, таких как MAGMA, прежде чем внедрять свое собственное.