Как выделить массив разделяемой памяти в CUDA Fortran? - PullRequest
Новая
       11

Как выделить массив разделяемой памяти в CUDA Fortran?

0 голосов
/ 09 декабря 2010

У меня проблемы с попыткой объявить массив общей памяти в ядре. Вот код, содержащий мое ядро: 

module my_kernels

  use cudafor
  implicit none

contains

  attributes(global) subroutine mykernel(N)

    ! Declare variables
    integer :: index
    integer, intent(in), value :: N
    real,shared,dimension(N) :: shared_array  

    ! Map threadID to index
    index = blockDim%x * (blockIdx%x-1) + threadIdx%x

    ! Set array element equal to index
    shared_array(index) = index

  end subroutine mykernel

end module my_kernels

А вот как я называю свое ядро: 

program cuda

  use my_kernels
  implicit none  

  ! Set number of threads
  integer :: N = 9

  ! Invoke kernel with 3 blocks of 3 threads
  call mykernel<<<N/3,3>>>(N)

end program cuda

Все это у меня в одном файле, test.cuf. Когда я пытаюсь скомпилировать test.cuf с pgf90, я получаю эту ошибку: 

PGF90-S-0000-Internal compiler error. unexpected runtime function call       0 (test.cuf: 34)
PGF90-S-0000-Internal compiler error. unsupported procedure     349 (test.cuf: 34)
  0 inform,   0 warnings,   2 severes, 0 fatal for mykernel
/tmp/pgcudaforw5MgcaFALD9p.gpu(19): error: a value of type "int" cannot be assigned to an entity of type "float *"

/tmp/pgcudaforw5MgcaFALD9p.gpu(22): error: expected an expression

2 errors detected in the compilation of "/tmp/pgnvdl7MgHLY1VOV5.nv0".
PGF90-F-0000-Internal compiler error. pgnvd job exited with nonzero status code       0 (test.cuf: 34)
PGF90/x86-64 Linux 10.8-0: compilation aborted

В этом случае строка 34 относится к end subroutine mykernel. Ошибка компилятора не очень полезна, мне потребовалось некоторое время, чтобы выяснить, что проблема была в общем массиве (я использую этот код в качестве простого примера).

Когда я заменяю 'N' на '9' в объявлении общего массива, так что real,shared,dimension(N) :: shared_array заменяется на real,shared,dimension(9) :: shared_array, ошибка исчезает.

У меня вопрос: почему возникает эта ошибка, и как мне задать размер общего массива с помощью переменной (если это действительно возможно)?

Ответы [ 2 ]

1 голос
/ 02 июня 2012

Вы можете иметь более одного массива совместно используемой памяти, но их размер должен быть известен во время компиляции.Обычно массивы совместно используемой памяти должны иметь фиксированный размер, случай, когда вы можете передать размер в байтах во время выполнения, является исключительным.Я думаю, это все из-за ограничения общей памяти в SM (потоковый мультипроцессор).По моему опыту разработки на CUDA C и CUDA fortran лучше, чтобы все эти параметры были «фиксированы», а затем заставляло ядро ​​повторять работу столько раз, сколько необходимо для охвата всех входных данных, таким образом, мне легче контролировать все параметры,влияет на занятость (насколько хорошо вы используете все физические ресурсы в GPU).

1 голос
/ 10 декабря 2010

Измените «dimension (N)» на «dimension (*)», а затем передайте размер общего массива (в байтах) в качестве третьего аргумента запуска вашего ядра.

Надеюсь, это поможет,

Mat 

% cat test.cuf 
module my_kernels

  use cudafor
  implicit none

  real, dimension(:), allocatable,device :: Ad
  real, dimension(:),allocatable :: Ah

contains

  attributes(global) subroutine mykernel(N)

    ! Declare variables
    integer :: index
    integer, intent(IN), value :: N
    real,shared,dimension(*) :: shared_array  

    ! Map threadID to index
    index = blockDim%x * (blockIdx%x-1) + threadIdx%x

    ! Set array element equal to index
    shared_array(index) = index

    Ad(index) = index

  end subroutine mykernel

end module my_kernels


program cuda

  use my_kernels
  implicit none  

  ! Set number of threads
  integer :: N = 9

   allocate(Ad(N), Ah(N))

  ! Invoke kernel with 3 blocks of 3 threads
  call mykernel<<<N/3,3,N*4>>>(N)

  Ah=Ad
  print *, Ah

end program cuda

% pgf90 test.cuf -V10.9 ; a.out
    1.000000        2.000000        3.000000        4.000000     
    5.000000        6.000000        7.000000        8.000000     
    9.000000
...