• Add options that prevent generation of debug code into the HIP C++ kernels
• Add options to disable prettification of generated HIP C++ code and formatted
  Fortran source

Background:

• Some applications might want to use GPUFORT as a preprocessing step in an automated
  compilation process. Output formatting is not needed in this case.

The following code will generate a parsing error:

program test1 integer, parameter :: ims = 12,ime=16,jms=1,jme=10 REAL, DIMENSION(ims:ime, jms:jme), INTENT(OUT) :: & ! Does not like a comment here HEAT2D ! What about here end program test1

• Translation of CUDA Fortran global subroutines is currently disabled

• Planned code generation:

Update(05/26/2021):

An initial implementation was pushed to the main branch and an example vector-add-kernel
was added as demonstator. A current limitation is that kernels can only call device and host-device functions
with scalar parameters.

Add:

• amax1

• amin1

• nint

• ...

• Labeled DO-WHILE loop

• parse interfaces on module level and give subprograms an interface name

Fix:

• max seems not to be implemented correctly

• There is a number of libraries we can take look on such as rapidjson, orjson.

• Associate generated HIP C++ files with modules / programs to allow inclusion of device routines or ...
• Inline everything straightforwardly?

The latter option seems to be the preferred one as it would also allow
to handle gang/worker/vector accelerator routines. Would require indexer to record whole function definitions.

The collapse directive is not preserved when converting from OpenACC to OpenMP.

• as a first step, we will allow the user to predefine a macro
• as second step, we will discover macros directly in the source code

• GPUFORT has seen dramatic changes in the past to enable OpenACC
• We need to reassess the CUDA Fortran porting capabilities of the code
• We further need to adopt the new HIPFORT Fortran interfaces

• Allows to convert only compute constructs and manually tune them while keeping all the other
  directives as directives in the code.
• The other directives can then be converted automatically by a preprocessor when
  compiling (add a separate option that excludes compute constructs).
• This would increase the maintainability of GPUFORT-translated applications.

Convert code such like:

program main; implicit none; end program main

program main
implicit none
end program main

• Split launchers and kernels into separate files?
• Reorder kernels and launchers?
• Scan for __global__ and extern "C"?

Issue:

Fortran derived types can be based on other types.
GPUFORT does not support this yet.

Background/Example:

type mytype1
   real :: a
end type

type, extend(mytype1) :: mytype2
   real :: b
end type

mytype2 inherited member field a from mytype1.

Assessment:

Issue has low priority as we did not see such code yet in any HPC
application that we tried to port.

• Double/custom precision numbers have a suffix that indicates the desired number format.
• Numbers without a suffix need to be translated to the default real (=float if not redefined via compiler option).

!$acc declare is not handled yet.

• "A declare directive is used in the declaration section of a Fortran subroutine,
  function, or module."

• "In a Fortran module declaration section, only create, copyin, device_resident, and
  2084 link clauses are allowed."

For more details on the directive, see "2.13. Declare Directive" in the specification

Implementation approach:

• Host functions & subroutines:

  • Track function/subroutine begin and declarations in scanner tree & indexer
  • Consider !$acc declare in indexer
  • Track return statements in scanner tree as child node of function/procedure (begin)
  • Track function/subroutine end in scanner tree as child node of function/procedure (begin)
  • Emit unstructured enter data routine call or equivalent shortcut after declaration section of function/subroutine
  • Emit unstructured exit data routine call or equivalent shortcut before function/subroutine end or return statement

• Modules / Program

  • Track module and declarations in scanner tree & indexer
  • Consider !$acc declare in indexer
  • Assumed-size arrays
    * [x] Track allocate statements in in scanner tree
    * [ ] Emit unstructured enter data routine call or equivalent shortcut after allocate
  • Explicit-size arrays
    • ??? - not totally sure where to put the allocation
      • if no save keyword:
        • body of function/subroutine/program that uses the module
      • else if save keyword and not already present in global scope (scope 0):
        • body of function/subroutine/program that uses the module
    • Emit unstructured enter data routine call or equivalent shortcut where appropriate

Accelerator routines use "acc-declared" variables

Some thoughts:

• Device routines can use variables that are mapped to the device via acc declare
  statements.
• When we generate HIP code, we need to pass the respective device pointers as kernel parameters.
• Hence, we would need to add a use statement with the respective module to the caller's (assuming loopnest) parent so that it is available in this scope. In case of gang, worker, vector accelerator routines, this must be done recursively.

In addition to the simple examples, I have previously provided. Several complex subroutines are also failing to parse properly. It was not immediately clear to me the reason for these failures. I have attached several examples. Note that despite the *.txt file extension, these are Fortran source files.

adv.txt
misclibs.txt
pdef.txt
solve1.txt
solve2.txt
soundcb.txt
sounde.txt

• Cannot specify new log directory via config when default location is inaccessible
• Conflicts when multiple users specify the same log file

It appears that the use of the pwr function ()**3 in parameter statements causes a parsing error. I have attached an example Fortran file.

module_mp_nssl_2mom.txt

When a Fortran line continuation character (&) is included in either an expression or an !$acc directive the parser fails to recognize the source code. See the example code snippet below

     ```

!$acc data copyin(a,b) copyout(c_gpu)
!$acc parallel loop collapse(2) &
!$acc reduction(+:tmp)
do j=1,colsB
do i=1,rowsA
tmp = 0.0
!$acc loop vector reduction(+:tmp)
do k=1,rowsB
tmp = tmp &
+ a(i,k) * b(k,j)
enddo
c_gpu(i,j) = tmp
enddo
enddo
!$acc end parallel
!$acc end data

The parser does not recognize the follow inlined if-statement:

if( tha(i,j,k).gt. 1.0 ) thrad = -1.0/(12.0*3600.0)

Background information

Single-line if statements need to be unrolled when the translated body
part will contain multiple statements that are subject to the condition part.

Steps to reproduce / example

We have the following statement that contains two CUDA Fortran device variables x_d, y_d.

if ( <condition> ) deallocate(x_d,y_d)

Expected outcome:

This should be converted to HIP as follows:

if ( <condition> ) then
  call hipCheck(hipFree(x_d))
  call hipCheck(hipFree(y_d))
else

Actual outcome

if ( <condition> )   call hipCheck(hipFree(x_d))          ! NOTE: end of single-line if statement
call hipCheck(hipFree(y_d))                                        ! NOTE: next statement

The second statement is not subject to the condition anymore, which is a bug.

GPUFORT can translate the following expressions as
they do not imply that an additional allocation and deallocation
must be generated at the begin and end of the scope.

real, device, allocatable :: x_d(:), y_d(:)

For the same reason, the translation of fixed-size array declarations is not
completely supported.

Hence,

 real, device :: x_d(N), y_d(N)

will simply be translated to

real, device, dimension(:) :: x_d
real, device, dimension(:) :: y_d

Currently, no corresponding hipMalloc and hipFree calls are generated into the code.

TBA

@gmarkomanolis;@jungwonkim;@sael9740;@sael9740;@gjost-git

We are going to enforce two factor authentication in (https://github.com/ROCmSoftwarePlatform/) organization on 29th April, 2022 .
Since we identified you as outside collaborator for ROCmSoftwarePlatform organization, you need to enable two factor authentication in your github account else you shall be removed from the organization after the enforcement.
Please skip if already done.

To set up two factor authentication, please go through the steps in below link:

https://docs.github.com/en/free-pro-team@latest/github/authenticating-to-github/configuring-two-factor-authentication

Please email "[email protected]" for queries

Current parser is too general and does not take into account that Fortran
statements such as end module mymodule or end do mylabel must be either
terminated with ; or a linebreak.

Examples (wanted behaviour):

Program Hello ! valid

Program Hello; ! valid

Program; Hello ! invalid

Program Hello; Print *, "Hello World" ! valid

Program Hello Print *, "Hello World" ! invalid

Program Hello
   Print *, "Hello World" ! valid

• In extracted kernels / device subroutines, we sometimes need to translate size(<array>[,<dim>]), lbound(<array>,<dim>)
  and ubound(<array>,<dim>) intrinsic calls,

Implementation status for different types of <array> and <dim>:

• <array> is identifier and <dim> is integer literal:
  • size(<array>,<dim>) -> <array>_n<dim>
  • lbound(<array>,<dim>) -> <array>_lb<dim>
  • ubound(<array>,<dim>) -> <array>_lb<dim> + <array>_n<dim>

NOTE: Above, the <array>_lb<dim> and <array>_n<dim> are already arguments of
the extracted routines.

• <array> is identifier and <dim> is identifier | arithmetic expression

  • In this case, we need to generate a query function that is parameterized by array rank.

• Other cases where <array> is not an identifier and / or <dim> is identifier are not supported

• Support acc routine

After the most recent update that addressed issue #34 a large number of index errors are now occurring in the spit_fortran_line python function. I have attached a Fortran code that reproduces this error.

Thanks,
John

writeout.txt

In CUDA Fortran, qualifiers of a variable can be added after the initial declaration.

Example:

real :: a(4)

#ifdef __CUDA
attributes(device) :: a
#endif

This must be taken into account by the indexer.

The develop-acc branch version of gpufort fails to parse the following file:

program test1 logical, parameter :: dosfcflx = .true. end program test1

This will allow us to call the CPU variant of a kernel from C++ code, which, in turn, will allow us to write/generate unit tests in C++.

The develop-acc branch fails to properly parse the following bit of code:

program test1 real, dimension(:,:,:,:), pointer :: tauaer_sw=>null() end program test1