Perform managed and device arrays on different streams

Accelerated Computing HPC Compilers Legacy PGI Compilers
1

I tried the example code in Account Login | PGI but got an error: Bus error (core dumped).
The code is here:

module mCuda
  integer :: num_cm
  type tCM
    integer, allocatable, device :: fine(:)
    real, allocatable, device :: mat_matrix(:,:,:)
    real, allocatable, device :: src_matrix(:,:,:)
  end type
  type(tCM), allocatable, managed :: cm_list(:)
end module

program test
use mCuda
use cudafor
integer(kind=cuda_stream_kind) :: str_mng, str_dev
integer i
num_cm=10
istat = cudaStreamCreate(str_mng)
istat = cudaStreamCreate(str_dev)
allocate(cm_list(num_cm))
istat = cudaforSetDefaultStream(cm_list, str_mng)
do i=1, num_cm
  allocate(cm_list(i)%fine(3))
  istat = cudaforSetDefaultStream(cm_list(i)%fine, str_dev)
  cm_list(i)%fine=10
enddo
end

Thanks for your attention!

2

What type of system are you running on? Usually when you use managed memory and get a bus error, it means the host and device is accessing managed memory at the same time.

3

The hardware is Tesla K80, and the compiler platform is PrgEnv-pgi/16.1. Thanks!

4

Brent’s intuition is correct. For older cards such as a K80, if managed memory is accessed from both the host and device at the same time, you’ll get a bus error. The same issue does not occur on new P100 or V100 devices.

Here, the problem line is “cm_list(i)%fine=10” since you’re accessing managed memory on the host via “cm_list(i)”, but also accessing it on the device when assigning to device array "fine.

The simple solution is to make the “device” arrays “managed”. This will have the effect of updating the array on the host, but have them accessible on the device.

-Mat

5

Thanks guys for the clarification!