Расчет PI с помощью Fortran & CUDA

Question

Я пытаюсь сделать простую программу в Фортран компилятор PGI. Эта простая программа будет использовать графическую карту для вычисления числа пи с помощью алгоритма "дартс". После долгой борьбы с этой программой, я наконец-то получил ее по большей части. Тем не менее, я в настоящее время застрял на передаче результатов должным образом. Должен сказать, что это довольно сложная программа для отладки, поскольку я больше не могу вставлять какие-либо операторы печати в подпрограмму. Эта программа в настоящее время возвращает все нули. Я не совсем уверен, что происходит, но у меня есть две идеи. Оба из которых я не уверен, как исправить:

1. Ядро CUDA не работает как-то?
2. Я не конвертирую значения должным образом? pi_parts = pi_parts_d

Ну, это состояние моей текущей программы. Все переменные с _d в конце обозначаем подготовленную CUDA память устройства, где все остальные переменные (за исключением ядра CUDA) являются типичными подготовленными переменными Fortran CPU. Теперь есть некоторые заявления для печати, которые я прокомментировал, что я уже попробовал с земли Фортрана. Эти команды должны были проверить, действительно ли я генерировал случайные числа правильно. Что касается метода CUDA, я в настоящее время закомментировал расчеты и заменил z статически равно 1 просто чтобы увидеть, как что-то происходит.

module calcPi
contains
    attributes(global) subroutine pi_darts(x, y, results, N)
        use cudafor
        implicit none
        integer :: id
        integer, value :: N
        real, dimension(N) :: x, y, results
        real :: z

        id = (blockIdx%x-1)*blockDim%x + threadIdx%x

        if (id .lt. N) then
            ! SQRT NOT NEEDED, SQRT(1) === 1
            ! Anything above and below 1 would stay the same even with the applied
            ! sqrt function. Therefore using the sqrt function wastes GPU time.
            z = 1.0
            !z = x(id)*x(id)+y(id)*y(id)
            !if (z .lt. 1.0) then
            !   z = 1.0
            !else
            !   z = 0.0
            !endif
            results(id) = z
        endif
    end subroutine pi_darts
end module calcPi

program final_project
    use calcPi
    use cudafor
    implicit none
    integer, parameter :: N = 400
    integer :: i
    real, dimension(N) :: x, y, pi_parts
    real, dimension(N), device :: x_d, y_d, pi_parts_d
    type(dim3) :: grid, tBlock

    ! Initialize the random number generaters seed
    call random_seed()

    ! Make sure we initialize the parts with 0
    pi_parts = 0

    ! Prepare the random numbers (These cannot be generated from inside the
    ! cuda kernel)
    call random_number(x)
    call random_number(y)

    !write(*,*) x, y

    ! Convert the random numbers into graphics card memory land!
    x_d = x
    y_d = y
    pi_parts_d = pi_parts

    ! For the cuda kernel
    tBlock = dim3(256,1,1)
    grid = dim3((N/tBlock%x)+1,1,1)

    ! Start the cuda kernel
    call pi_darts<<<grid, tblock>>>(x_d, y_d, pi_parts_d, N)

    ! Transform the results into CPU Memory
    pi_parts = pi_parts_d
    write(*,*) pi_parts

    write(*,*) 'PI: ', 4.0*sum(pi_parts)/N
end program final_project

РЕДАКТИРОВАТЬ К КОДУ: Изменены различные строки, чтобы отразить исправления, упомянутые: Robert Crovella, Текущий статус: ошибка обнаружена cuda-memcheck выявление: Program hit error 8 on CUDA API call to cudaLaunch на моей машине.

Если есть какой-либо метод, который я могу использовать для тестирования этой программы, пожалуйста, дайте мне знать. Я бросаю дротики и смотрю, где они приземляются для моего текущего стиля отладки с CUDA. Не самый идеальный, но это придется делать, пока я не найду другой путь.

Пусть Боги Фортрана помилуют мою душу в этот темный час.

Answer 1

Когда я компилирую и запускаю вашу программу, у меня возникает ошибка. Это связано с последним параметром, который вы передаете ядру (N_d):

call pi_darts<<<grid, tblock>>>(x_d, y_d, pi_parts_d, N_d)

поскольку N это скалярное количество, ядро ​​ожидает использовать его напрямую, а не как указатель. Поэтому, когда вы передаете указатель на данные устройства (N_d), процесс настройки ядра генерирует ошибку сегмента (в коде хоста!), когда он пытается получить доступ к значению N, который должен быть передан непосредственно как:

call pi_darts<<<grid, tblock>>>(x_d, y_d, pi_parts_d, N)

Когда я внесу это изменение в код, который вы опубликовали, я получу фактический вывод на печать (вместо ошибки сегмента), который представляет собой массив единиц и нулей (256 единиц, за которыми следуют 144 нуля), всего N=400 значений), за которым следует вычисленное значение PI (которое в данном случае составляет 2,56 (4*256/400), так как вы сделали ядро ​​в основном фиктивным ядром).

Эта строка кода также, вероятно, не делает то, что вы хотите:

grid = dim3(N/tBlock%x,1,1)

С N =400 и tBlock%x = 256 (из предыдущих строк кода), результат вычисления равен 1 (т.е. grid заканчивается в (1,1,1) что составляет один поток). Но вы действительно хотите запустить 2 потоковых блока, чтобы охватить весь диапазон вашего набора данных (N =400 элементов). Есть несколько способов исправить это, но для простоты давайте всегда добавим 1 к вычислению:

grid = dim3((N/tBlock%x)+1,1,1)

При этих обстоятельствах, когда мы запускаем сетки, которые больше (с точки зрения общего количества потоков), чем наш размер набора данных (512 потоков, но только 400 элементов данных в этом примере), обычно ставят проверку потока в начале нашего ядра (в этот случай, после инициализации id), чтобы предотвратить доступ за границу, например:

if (id .lt. N) then

(и соответствующий endif в самом конце кода ядра) Таким образом, только потоки, которые соответствуют действительным действительным данным, могут выполнять любую работу.

С учетом вышеуказанных изменений ваш код должен быть по существу функциональным, и вы должны иметь возможность вернуть код ядра в правильные операторы и начать получать оценку PI.

Обратите внимание, что вы можете проверить API CUDA на наличие кодов возврата ошибок, а также запустить код с cuda-memcheck чтобы получить представление о том, делает ли ядро ​​доступ за пределами допустимого. Однако ни один из них не помог бы с этой конкретной ошибкой сегмента.