Register pressure: Fortran

[a-v-medvedev]

In the "Register pressure" article you show nice compiler flags that help us to get the register and occupancy information on the compiling stage. The options are:
"-Rpass-analyze=kernel-resource-usage"
"--save-temps"

That is great for C++ code, but I can't find any equivalent for AMD fortran compiler. Is there any way to get this information with amdflang?

--Alexey

[afanfa]

Hi @a-v-medvedev ,
are you asking how you can get register and occupancy info for CPU only code from the amdflang compiler or for GPU code offloaded via OpenMP using the amdflang compiler?

In the latter case, try setting export LIBOMPTARGET_KERNEL_TRACE=1, you should get information about register usage when you run the code.

[a-v-medvedev]

In the latter case, try setting export LIBOMPTARGET_KERNEL_TRACE=1, you should get information about register usage when you run the code.

Thanks for your answer,
yes I am definitely asking about OpenMP GPU offloading case. Thanks for a great hint, I was not aware of this environment variable.

But originally I was wondering if I can get this information for Fortran code at compile time, just the similar way as it is done with "-Rpass-analyze=kernel-resource-usage", "--save-temps" for C++.

--Alexey

[jychang48]

@a-v-medvedev try these options:

-mllvm -amdgpu-dump-hsa-metadata

[a-v-medvedev]

@jychang48 @afanfa Nope, for fortran this doesn't work...

Moreover, LIBOMPTARGET_KERNEL_TRACE also doesn't work for Cray Fortran (from CCE 15), but CRAY_ACC_DEBUG=3 produces just something:

ACC:       blocks:  28
ACC:      threads:  128
...
ACC:    getting function wrapper$wrapper_mod_$ck_L124_1_cce$noloop$form
ACC:       stats threads=1024 threadblocks per cu=8 shared=0 total shared=0
ACC:       prefer equal shared memory and L1 cache 
ACC:     kernel information
ACC:               num registers :      128
ACC:        max theads per block :     1024
ACC:                 shared size :        0 bytes
ACC:                  const size :        0 bytes
ACC:                  local size :     1424 bytes

From here, I can deduce that: "threads: 128" -> 2 wavefront per block; "threadblocks per cu=8" -> 16 wavefronts per cu; and we know that the maximum is 32 for MI250x; also we understand, that the number of registers is 128. Better than nothing.

[jychang48]

@a-v-medvedev which fortran compiler are you using specifically? If you're using /opt/rocm-5.7.0/bin/amdflang on MI200's you need these full options:

amdflang -fopenmp -mllvm  -amdgpu-dump-hsa-metadata --offload-arch=gfx90a -D__OFFLOAD_ARCH_gfx90a__ -O3  ...

You should immediately see output looking something like this:

AMDGPU HSA Metadata:
---
amdhsa.kernels:
  - .agpr_count:     0
    .args:
      - .name:           Arg_n_0.arg
        .offset:         0
        .size:           8
        .value_kind:     by_value
      - .address_space:  generic
        .name:           'Arg_output$sd'
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  generic
        .name:           'Arg_output$p'
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 56
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           __nv_MAIN__TARGET_F1L37_1_
    .private_segment_fixed_size: 0
    .sgpr_count:     36
    .sgpr_spill_count: 0
    .symbol:         __nv_MAIN__TARGET_F1L37_1_.kd
    .uses_dynamic_stack: false
    .vgpr_count:     22
    .vgpr_spill_count: 0
    .wavefront_size: 64
amdhsa.target:   amdgcn-amd-amdhsa--gfx90a
amdhsa.version:
  - 1
  - 2
...

The environment variable LIBOMPTARGET_KERNEL_TRACE is applicable only to AMD's fortran compiler. Whatever information you extract from CRAY_ACC_DEBUG will not necessarily be true for a code compiled with amdflang. Some recommendations:

1. If you have access to the Cray Fortran, please use it for your GPU offloading work.
2. amdflang is based on the deprecated classic flang, you likely won't get good performance out of it from more complex kernels. AMD compiler engineers are still working on a new and performant flang but there is no release ETA at the moment. In the meantime you could try using one of the AOMP releases which again is still based on classic flang but has newer builds: https://github.com/ROCm-Developer-Tools/aomp

[a-v-medvedev]

@jychang48 sorry for a delay in answer. We are using CCE 16 Fortran compiler now. In fact, the CRAY_ACC_DEBUG output is not providing the occupancy-related info. I do not think we are able to switch to amdflang easily in the compilation flow that we have, but we probably can do this for small reproducer codes.

Still lacking any way to get real observations on occupancy for selected kernels using Cray CCE Fortran + OpenMP target offloading. We can only use theoretical deducing of the occupancy figure based on number of registers and amount of LDS used, which is not convenient and error-prone.