Оптимизация CUDA FDTD Фортран - PullRequest
Новая
       5

Оптимизация CUDA FDTD Фортран

1 голос
/ 04 марта 2012

Я пытаюсь оптимизировать этот код FDTD с помощью CUDA Fortran. У меня есть три трехмерных куба матрицы с входом, выходом и стоимостью. 

attributes (global) subroutine kernel_h(k,num_cells_x,num_cells_y,num_cells_z,Hx,Hy,Hz,Ex,Ey,Ez,Cbdx,Cbdy,Cbdz)
    implicit none
    integer :: idx,idy
    integer,value :: k,num_cells_x,num_cells_y,num_cells_z
    real(kind=8), intent(in), dimension(1:num_cells_x,1:num_cells_y,1:num_cells_z) :: Ex, Ey, Ez
    real(kind=8), intent(inout), dimension(1:num_cells_x,1:num_cells_y,1:num_cells_z) :: Hx, Hy, Hz
    real(kind=8), intent(in), constant, dimension(1:num_cells_x,1:num_cells_y,1:num_cells_z) :: Cbdx,Cbdy,Cbdz
    idx = threadIdx%x + ((blockIdx%x-1) * blockDim%x)
    idy = threadIdx%y + ((blockIdx%y-1) * blockDim%y)
    do while (idx < num_cells_x)
        Hz(idx,idy,k) = Hz(idx,idy,k) + ((Ex(idx,idy+1,k)-Ex(idx,idy,k))*Cbdy(idx,idy,k) + (Ey(idx,idy,k)-Ey(idx+1,idy,k))*Cbdx(idx,idy,k))
        Hx(idx,idy,k) = Hx(idx,idy,k) + ((Ey(idx,idy,k+1)-Ey(idx,idy,k))*Cbdz(idx,idy,k) + (Ez(idx,idy,k)-Ez(idx,idy+1,k))*Cbdy(idx,idy,k))
        Hy(idx,idy,k) = Hy(idx,idy,k) + ((Ez(idx+1,idy,k)-Ez(idx,idy,k))*Cbdx(idx,idy,k) + (Ex(idx,idy,k)-Ex(idx,idy,k+1))*Cbdz(idx,idy,k))
        idx = idx + (blockDim%x * gridDim%x)
        idy = idy + (blockDim%y * gridDim%y)
    end do
end subroutine kernel_h

и мой запуск ядра: 

bdim=dim3(16,16,1)
gdim=dim3((num_cells_x+(bdim%x-1))/bdim%x,(num_cells_y+(bdim%y-1))/bdim%y,1)
do k=1,num_cells_z
 call kernel_h<<<gdim,bdim>>>(k,num_cells_x,num_cells_y,num_cells_z,Hx_d,Hy_d,Hz_d,Ex_d,Ey_d,Ez_d,Cbdx_d,Cbdy_d,Cbdz_d)
end do

Мои вопросы: почему я не могу загрузить матрицу размером более 100x100x100? Если я пытаюсь, я получаю ошибку запуска ядра. И могу ли я улучшить производительность моего кода? Я думаю, что это можно было бы написать лучше.

1 Ответ

3 голосов
/ 08 марта 2012

Полагаю, вы выходите за пределы.

Рассмотрим объем 10x10x10 (x, y, z). В этом случае вы запустите один блок из потоков 16x16. Эти нити получат доступ к срезу 17x17 (так как радиус трафарета равен 1), который явно выйдет за пределы. Вам нужно будет отключить те потоки, которые будут выходить за пределы, а также отключить те потоки, которые достигнут границы, чтобы применить их трафарет.

Рассмотрите пример FDTD3D в CUDA SDK. Конечно, он написан на C, но иллюстрирует, как справиться с этой проблемой, а также показывает, как использовать разделяемую память, чтобы иметь намного более эффективную реализацию.

...