Fortran+OpenACCのsubroutine内可変長配列

なんの話か

Fortran/OpenACCのループ内サブルーチンで可変長配列を用いた場合、NVIDIA GPU上ではどこのメモリにどのタイミングで確保されるか調べてみました。
コンパイル時に配列長が決定していないためレジスタに作られることはないだろうとの予想のもと、カーネル内mallocを呼んでいるのかなーと言う気持ちで調べていきます。

実験コード

調査に用いたコードはこちらです。
inline展開されると困るので、2ファイルに分けて書いています。

▼sub.f90

subroutine sub_routine(a, array_size)
  !$acc routine seq
  real(8) :: a(:)
  integer :: array_size

  real(8), dimension(array_size) :: array
  integer :: i

  do i = 1, array_size
    a(i) = array(i)
  end do
end subroutine

▼main.f90

program main
  interface
    subroutine sub_routine(a, array_size)
      !$acc routine seq
      real(8) :: a(:)
      integer :: array_size
    end subroutine
  end interface


  integer :: array_size = 10000
  integer :: n = 100000

  real(8), dimension(:), allocatable :: a

  allocate(a(n))

  !$acc kernels copy(a(1:n)), copyin(array_size)
  do i = 1, n
  call sub_routine(a, array_size)
  enddo
  !$acc end kernels

end program

（Makefile等はこちら enp1s0/openacc-subroutine - GitHub）

2ファイル用意したものの、実際にはsub.f90があれば調査できます。
コンパイルしたsub.o内に含まれるPTXをcuobjdumpで覗きます。

!...

$L__tmp0:
.loc	1 1 1
ld.u64 %rd2, [%rd22+56];
ld.u64 %rd3, [%rd22+80];
add.s64 %rd24, %rd2, %rd3;
add.s64 %rd44, %rd24, -1;
ld.u32 %r1, [%rd21];
mul.wide.s32 %rd25, %r1, 8;
{
	.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd25;
.param .b64 retval0;
call.uni (retval0),
malloc,
(
param0
);
ld.param.b64 %rd5, [retval0+0];
}
	setp.ne.s64 %p1, %rd5, 0;
@%p1 bra $L__BB0_4;

!...

予想通りmallocが呼ばれていることが確認できます。
つまり可変長配列の実体はDeviceメモリ上で、カーネル実行の最中に確保されるようです。
では、可変長配列をサイズ決め打ちの固定長にするとどうなるかと言うと、

サイズが小さいときはレジスタ
サイズが多いときは同様にカーネル内malloc

となることが確認できます。

おわり

カーネル内mallocというものを知った時誰が使うのだろうと思っていましたが、OpenACC/Fortranの実装で使われていたのですね。