Изменить значения нескольких ячеек в CUDA

Это должно быть просто, хотя я не могу найти ответ. Я пишу программу, которая должна вычислять состояния клеточных автоматов, и чтобы понять, как работает CUDA, я сначала попытался написать очень простую программу. Он принимает матрицу, и каждый поток должен увеличивать значение в своей ячейке и в ячейках, которые выше и ниже этой ячейки. Итак, если я дам ему следующую матрицу:

[0 0 0 0 0 0 0]
[0 0 0 0 0 0 0]
[0 0 0 0 0 0 0]
[0 0 0 0 0 0 0]
[0 0 0 0 0 0 0]
[0 0 0 0 0 0 0]
[0 0 0 0 0 0 0]

Я ожидаю получить следующий результат:

[2 2 2 2 2 2 2]
[3 3 3 3 3 3 3]
[3 3 3 3 3 3 3]
[3 3 3 3 3 3 3]
[3 3 3 3 3 3 3]
[3 3 3 3 3 3 3]
[2 2 2 2 2 2 2]

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

[2 2 2 2 2 2 2]
[3 3 3 3 3 3 3]
[3 3 3 3 3 3 3]
[3 3 3 3 2 2 2]
[2 2 2 2 2 2 2]
[2 2 2 2 3 3 3]
[2 2 2 2 2 2 2]

И я не могу понять, почему в 4-й, 5-й и 6-й строке стоят значения 2 - должно быть 3, а не 2.
Вот мой код:

import numpy
import pycuda.autoinit
import pycuda.driver as cuda

from pycuda.compiler import SourceModule

w = 7

mod = SourceModule("""
        __global__ void diffusion(  int* result, int width, int height) {

            int xIndex = blockDim.x * blockIdx.x + threadIdx.x;
            int yIndex = blockDim.y * blockIdx.y + threadIdx.y;

            int flatIndex = xIndex + width * yIndex;
            int topIndex = xIndex + width * (yIndex - 1);
            int bottomIndex = xIndex + width * (yIndex + 1);

            int inc = 1;

            result[flatIndex] += inc;

            result[bottomIndex] += inc;

            result[topIndex] += inc;
        }

        """)

diff_func   = mod.get_function("diffusion")


def diffusion(res):

    height, width = numpy.int32(len(res)), numpy.int32(len(res[0]))

    diff_func(
        cuda.InOut(res),
        width,
        height,
        block=(w,w,1)
        )

def run(res, step):

    diffusion(res)
    print res

res   = numpy.array([[0 \
                        for _ in xrange(0, w)]\
                        for _ in xrange(0, w)], dtype='int32')

run(res, 0)

Еще одна интересная вещь: если я прокомментирую одну из следующих строк:

result[bottomIndex] += inc;
result[topIndex] += inc;

Все работает так, как ожидалось, и нет никаких неожиданных значений. Похоже, в некоторых случаях CUDA не может работать со значениями трех соседних ячеек в одном потоке.


cuda pycuda
person aga    schedule 10.02.2013    source источник

Ответы (1)


arrow_upward
2
arrow_downward

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

Решение состоит в том, чтобы либо использовать операции с атомарной памятью (дополнительную информацию см. в руководстве по программированию CUDA), либо использовать другой подход для обновления соседних ячеек (например, раскрасить сетку и обновить как окрашенные ячейки при отдельных проходах по сетке).

person talonmies    schedule 10.02.2013
comment
Я думал об этом, но не исследовал, потому что ячейки с неожиданными значениями всегда одинаковы - я могу вызвать свою программу десять раз подряд, и всегда будет 2 в строках с 4-й по 6-ю, а не со 2-й по 6-ю. 3-й или только в третьем например. Я думал, что если есть какая-то гонка за ресурсами, то должна быть какая-то непредсказуемость в результатах. В любом случае, спасибо за ответ, я подробнее рассмотрю гонки памяти в CUDA. :) - person aga; 10.02.2013