PyCUDA незаконный доступ к памяти curandState * - PullRequest
Новая
       28

PyCUDA незаконный доступ к памяти curandState *

0 голосов
/ 23 июня 2019

Я изучаю распространение инвазивных видов и пытаюсь генерировать случайные числа в ядре PyCUDA с помощью генератора случайных чисел XORWOW.Матрицы, которые мне нужно использовать в качестве входных данных для исследования, достаточно велики (до 8 000 x 8 000).

Кажется, что ошибка возникает внутри get_random_number при индексации curandState* генератора XORWOW,Код выполняется без ошибок на меньших матрицах и дает правильные результаты.Я использую свой код на 2 графических процессорах NVidia Tesla K20X.

Код ядра и настройка: 

kernel_code = '''
    #include <curand_kernel.h>
    #include <math.h>

    extern "C" {

    __device__ float get_random_number(curandState* global_state, int thread_id) {

        curandState local_state = global_state[thread_id];
        float num = curand_uniform(&local_state);
        global_state[thread_id] = local_state;
        return num;
    }

    __global__ void survival_of_the_fittest(float* grid_a, float* grid_b, curandState* global_state, int grid_size, float* survival_probabilities) {

        int x = threadIdx.x + blockIdx.x * blockDim.x;             // column index of cell
        int y = threadIdx.y + blockIdx.y * blockDim.y;             // row index of cell

        // make sure this cell is within bounds of grid
        if (x < grid_size && y < grid_size) {

            int thread_id = y * grid_size + x;                      // thread index
            grid_b[thread_id] = grid_a[thread_id];                  // copy current cell
            float num;

            // ignore cell if it is not already populated
            if (grid_a[thread_id] > 0.0) {

                num = get_random_number(global_state, thread_id);

                // agents in this cell die
                if (num < survival_probabilities[thread_id]) {
                    grid_b[thread_id] = 0.0;                        // cell dies
                    //printf("Cell (%d,%d) died (probability of death was %f)\\n", x, y, survival_probabilities[thread_id]);
                }
            }
        }
    }

mod = SourceModule(kernel_code, no_extern_c = True)
survival = mod.get_function('survival_of_the_fittest')

Настройка данных: 

matrix_size = 2000
block_dims = 32
grid_dims = (matrix_size + block_dims - 1) // block_dims

grid_a = gpuarray.to_gpu(np.ones((matrix_size,matrix_size)).astype(np.float32))
grid_b = gpuarray.to_gpu(np.zeros((matrix_size,matrix_size)).astype(np.float32))
generator = curandom.XORWOWRandomNumberGenerator()
grid_size = np.int32(matrix_size)
survival_probabilities = gpuarray.to_gpu(np.random.uniform(0,1,(matrix_size,matrix_size)))

Вызов ядра: 

survival(grid_a, grid_b, generator.state, grid_size, survival_probabilities, 
    grid = (grid_dims, grid_dims), block = (block_dims, block_dims, 1))

Я ожидаю, что смогу генерировать случайные числа в пределах диапазона (0,1) для матриц до (8000 x 8000), но выполнение моего кода на больших матрицах приводит к недопустимой ошибке доступа к памяти. 

pycuda._driver.LogicError: cuMemcpyDtoH failed: an illegal memory access was encountered
PyCUDA WARNING: a clean-up operation failed (dead context maybe?)
cuMemFree failed: an illegal memory access was encountered

Неправильно ли я индексирую curandState* в get_random_number? И если нет, то что еще может вызывать эту ошибку?

1 Ответ

2 голосов
/ 25 июня 2019

Проблема здесь заключается в разрыве между этим кодом , который определяет размер состояния, которое интерфейс PyCUDA curandom выделяет для своего внутреннего состояния, и этим кодом в вашем сообщении: 

matrix_size = 2000
block_dims = 32
grid_dims = (matrix_size + block_dims - 1) // block_dims

Похоже, вы предполагаете, что PyCUDA волшебным образом выделит достаточно состояния для любого размера блока и сетки, которые вы выберете в своем коде. Это маловероятно, особенно при больших размерах сетки. Вам либо нужно

  • Измените код так, чтобы он использовал те же размеры блоков и сеток, что и модуль curandom, используемый внутри для того генератора, который вы выбрали, или
  • Выделите и управляйте собственным пространством пустых состояний, чтобы у вас было достаточно выделенного состояния для обслуживания выбранных блоков и размеров сетки

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

...