1

I’m new to CUDA Fortran. Can someone please explain what is wrong with the following code? Why am I getting segmentation fault? What should I do to pass variables from different modules to a kernel called a subroutine as the code below?

case 1:

attributes(global) subroutine increment(a)
    implicit none
    real(kind=8),device::a(:)
    integer :: i
    i = threadIdx%x
    a(i)=a(i)+1
    if(a(i)/=2)then
        write(*,*)' Program Passed in Device'
    endif
end subroutine increment

module common
    implicit none
    real(kind=8):: a(2)
    real(kind=8),device:: a_d(2)
end module common
    
program incrementTestGPU
    use cudafor
    use common
    implicit none
    a=1
    a_d=a
    call prac    
end program incrementTestGPU

subroutine prac
    use common
    use cudafor
    call increment<<<1,2>>>(a_d)
    a=a_d
    if(any(a/=2))then
        write(*,*)' **** Program Failed **** '
    else
        write(*,*)' **** Program Passed **** '
    endif
end subroutine prac

Earlier I got segmentation fault, now the code is not printing the write statement in the device.

case 2:

module simpleOps_m
   contains
   attributes(global) subroutine increment(a)
   implicit none
   real(kind=8),device::a(2)
   integer :: i
   i = threadIdx%x
   a(i)=a(i)+1
   if(a(i)/=2)then
       print*,' Program Passed in Device',a(i)
   endif
end subroutine increment
end module simpleOps_m

module common
   implicit none
   real(kind=8):: a(2)
   real(kind=8),device:: a_d(2)
end module common
   
program incrementTestGPU
   use cudafor
   use common
   use practice
   implicit none
   a=1
   a_d=a
   call prac    
end program incrementTestGPU

module practice
   contains
   subroutine prac
       use simpleOps_m
       call increment<<<1,2>>>(a_d)
       a=a_d
       if(any(a/=2))then
           write(*,*)' **** Program Failed **** '
       else
           write(*,*)' **** Program Passed **** '
       endif
   end subroutine prac
end module practice```
Nethaji
  • 63
  • 8
  • Welcome, I suggest taking the [tour]. Please use tag [tag:fortran] for Fortran questions. – Vladimir F Героям слава Aug 15 '21 at 06:23
  • 2
    `increment` is an external subroutine (not in the module), but has an assumed-shape dummy argument. Did you mean for it to have an assumed-size argument (`a(*)` instead of `a(:)`). I'm not entirely convinced that CUDA Fortran plays well with assumed-shape dummies but to give it a reasonable chance you'll likely need to give `increment` an explicit interface in `prac`. – francescalus Aug 15 '21 at 09:08
  • Hi, @francescalus can you please explain clearly. I don't understand what you are saying. I tried the a(*), a(:),a(2) as dummy arguments, the program fails everytime. – Nethaji Aug 15 '21 at 10:05
  • Hi, @francescalus, can you please explain clearly. I don't understand what you are saying. I tried a(*), a(:), a(2) as dummy arguments, the program fails every time. What do you mean give increment an explicit interface in prac...? The program is working fine if I change the kernel to a normal subroutine which increments array a in a loop. How is calling a kernel subroutine different from calling a normal subroutine...? – Nethaji Aug 15 '21 at 10:14
  • 1
    What you have written, ignoring the cuda stuff, is not valid - if you use assumed shape dummy arguments you must have a interface in scope at the calling point. The best way to do this is to put the subroutine in a module and use it, in fact I would strongly recommend you put all subprograms in a module. For more details see https://stackoverflow.com/questions/42766530/procedure-with-assumed-shape-dummy-argument-must-have-an-explicit-interface and https://stackoverflow.com/questions/13293964/module-calling-an-external-procedure-with-implicit-interface – Ian Bush Aug 15 '21 at 10:53
  • 1
    When calling the kernel subroutine data has to be _transferred_ to the device (copying arguments, or referencing things placed there otherwise): simple references to host memory just aren't suitable. – francescalus Aug 15 '21 at 12:19
  • Of course I did the transfer to the device @francescalus before calling the kernel, before line `call prac` right...? Or do you mean something else...? – Nethaji Aug 15 '21 at 17:03
  • Have you tried with the subroutines in the module? – francescalus Aug 15 '21 at 17:33
  • Yes, @francescalus I tried with subroutines in the module in case 2 – Nethaji Aug 16 '21 at 04:42
  • How are you compiling (to ensure you have enabled character processing and output in kernels)? You have `a(i)/=2` as the condition for printing, but isn't `2` the value you expect to have in `a`? – francescalus Aug 16 '21 at 08:42
  • I got it @francescalus. Thank you for your help. If you don't mind please help me with this question https://stackoverflow.com/questions/68798704/problem-with-using-managed-module-variables-from-other-files-in-cuda-fortran – Nethaji Aug 16 '21 at 10:40

0 Answers0