GPU Parallelization

GPU Parallelization

MFC compiles GPU code via OpenACC and in the future OpenMP as well.

In order to swap between OpenACC and OpenMP, custom GPU macros are used that translate to equivalent OpenACC and OpenMP directives. FYPP is used to process the GPU macros.

OpenACC Quick start Guide

OpenACC API Documentation

---

Macro API Documentation

Note: Ordering is not guaranteed or stable, so use key-value pairing when using macros

Data Type Meanings

• Integer is a number
• Boolean is a pythonic boolean - Valid options: True or False
• String List is given as a comma separated list surrounding by brackets and inside quotations
  • Ex: '[hello, world, Fortran]'
• 2-level string list is given as a comma separated list of string lists surrounding by brackets and inside quotations
  • Ex: '[[hello, world], [Fortran, MFC]]' or '[[hello]]'

Data Flow

• Data on the GPU has a reference counter
• When data is referred to being allocated, it means that GPU memory is allocated if it is not already present in GPU memory. If a variable is already present, the reference counter is just incremented.
• When data is referred to being deallocated, it means that the reference counter is decremented. If the reference counter is zero, then the data is actually deallocated from GPU memory
• When data is referred to being attached, it means that the device pointer attaches to target if it not already attached. If pointer is already attached, then the attachment counter is just incremented
• When data is referred to being detached, it means that the attachment counter is decremented. If attachment counter is zero, then actually detached

---

Computation Macros

GPU_PARALLEL_LOOP – (Execute the following loop on the GPU in parallel)

Macro Invocation

In order to parallelize a loop, simply place two macro calls on either end of the loop:

$:$GPU_PARALLEL_LOOP(...)
   {code}
$:END_GPU_PARALLEL_LOOP()

This wraps the lines in code with parallelization calls to openACC or openMP, depending on environment and compiler settings.

Parameters

name	data type	Default Value	description
code	code	Required	Region of code where the GPU parallelizes loops
collapse	integer	None	Number of loops to combine into 1 loop
parallelism	string list	'[gang,vector]'	Parallelism granularity to use for this loop
default	string	'present'	Implicit assumptions compiler should make
private	string list	None	Variables that are private to each iteration/thread
firstprivate	string list	None	Initialized variables that are private to each iteration/thread
reduction	2-level string list	None	Variables unique to each iteration and reduced at the end
reductionOp	string list	None	Operator that each list of reduction will reduce with
copy	string list	None	Allocates and copies data to GPU on entrance, then deallocated and copies to CPU on exit
copyin	string list	None	Allocates and copies data to GPU on entrance and then deallocated on exit
copyinReadOnly	string list	None	Allocates and copies readonly data to GPU and then deallocated on exit
copyout	string list	None	Allocates data on GPU on entrance and then deallocates and copies to CPU on exit
create	string list	None	Allocates data on GPU on entrance and then deallocates on exit
no_create	string list	None	Use data in CPU memory unless data is already in GPU memory (OpenACC only)
present	string list	None	Data that must be present in GPU memory. Increment counter on entrance, decrement on exit
deviceptr	string list	None	Pointer variables that are already allocated on GPU memory
attach	string list	None	Attaches device pointer to device targets on entrance, then detach on exit
extraAccArgs	string	None	String of any extra arguments added to the OpenACC directive
extraOmpArgs	string	None	String of any extra arguments added to the OpenMP directive

Parameter Restrictions

name	Restricted range
collapse	Must be greater than 1
parallelism	Valid elements: 'gang', 'worker', 'vector', 'seq'
default	'present' or 'none'

Additional information

• default present means that the any non-scalar data in assumed to be present on the GPU
• default none means that the compiler should not implicitly determine the data attributes for any variable
• reduction and reductionOp must match in length
• With reduction='[[sum1, sum2], [largest]]' and reductionOp='[+, max]', sum1 and sum2 will be the sum of sum1/sum2 in each loop iteration, and largest will the maximum value of largest all the loop iterations
• A reduction implies a copy, so it does not need to be added for both

Example

#:call GPU_PARALLEL_LOOP(collapse=3, private='[tmp, r]', reduction='[[vol, avg], [max_val]]', reductionOp='[+, MAX]')
  {code}
  ...
#:endcall GPU_PARALLEL_LOOP
#:call GPU_PARALLEL_LOOP(collapse=2, private='[sum_holder]', copyin='[starting_sum]', copyout='[eigenval,C]')
   {code}
  ...
#:endcall GPU_PARALLEL_LOOP

GPU_LOOP – (Execute loop on GPU)

Macro Invocation

Uses FYPP eval directive using $:

$:GPU_LOOP(...)

Parameters

name	data type	Default Value	description
collapse	integer	None	Number of loops to combine into 1 loop
parallelism	string list	None	Parallelism granularity to use for this loop
data_dependency	string	None	'independent'-> assert loop iterations are independent, 'auto->let compiler analyze dependencies
private	string list	None	Variables that are private to each iteration/thread
reduction	2-level string list	None	Variables unique to each iteration and reduced at the end
reductionOp	string list	None	Operator that each list of reduction will reduce with
extraAccArgs	string	None	String of any extra arguments added to the OpenACC directive
extraOmpArgs	string	None	String of any extra arguments added to the OpenMP directive

Parameter Restrictions

name	Restricted range
collapse	Must be greater than 1
parallelism	Valid elements: 'gang', 'worker', 'vector', 'seq'
data_dependency	'auto' or 'independent'

Additional information

• Loop parallelism is most commonly '[seq]'
• reduction and reductionOp must match in length
• With reduction='[[sum1, sum2], [largest]]' and reductionOp='[+, max]', sum1 and sum2 will be the sum of sum1/sum2 in each loop iteration, and largest will the maximum value of largest all the loop iterations

Example

$:GPU_LOOP(parallelism='[seq]')
$:GPU_LOOP(collapse=3, parallelism='[seq]',private='[tmp, r]')

GPU_PARALLEL – (Execute the following on the GPU in parallel)

Macro Invocation

Uses FYPP call directive using #:call

#:call GPU_PARALLEL(...)
   {code}
#:endcall GPU_PARALLEL 

Parameters

name	data type	Default Value	description
code	code	Required	Region of code where a kernel is launched on the GPU
default	string	'present'	Implicit assumptions compiler should make
private	string list	None	Variables that are private to each iteration/thread
firstprivate	string list	None	Initialized variables that are private to each iteration/thread
reduction	2-level string list	None	Variables unique to each iteration and reduced at the end
reductionOp	string list	None	Operator that each list of reduction will reduce with
copy	string list	None	Allocates and copies data to GPU on entrance, then deallocated and copies to CPU on exit
copyin	string list	None	Allocates and copies data to GPU on entrance and then deallocated on exit
copyinReadOnly	string list	None	Allocates and copies readonly data to GPU and then deallocated on exit
copyout	string list	None	Allocates data on GPU on entrance and then deallocates and copies to CPU on exit
create	string list	None	Allocates data on GPU on entrance and then deallocates on exit
no_create	string list	None	Use data in CPU memory unless data is already in GPU memory (OpenACC only)
present	string list	None	Data that must be present in GPU memory. Increment counter on entrance, decrement on exit
deviceptr	string list	None	Pointer variables that are already allocated on GPU memory
attach	string list	None	Attaches device pointer to device targets on entrance, then detach on exit
extraAccArgs	string	None	String of any extra arguments added to the OpenACC directive
extraOmpArgs	string	None	String of any extra arguments added to the OpenMP directive

Parameter Restrictions

name	Restricted range
default	'present' or 'none'

Additional information

• default present means that the any non-scalar data in assumed to be present on the GPU
• default none means that the compiler should not implicitly determine the data attributes for any variable
• reduction and reductionOp must match in length
• With reduction='[[sum1, sum2], [largest]]' and reductionOp='[+, max]', sum1 and sum2 will be the sum of sum1/sum2 in each loop iteration, and largest will the maximum value of largest all the loop iterations
• A reduction implies a copy, so it does not need to be added for both

Example

#:call GPU_PARALLEL()
     {code}
     ...
#:endcall GPU_PARALLEL
#:call GPU_PARALLEL(create='[pixel_arr]', copyin='[initial_index]')
     {code}
     ...
#:endcall

---

Data Control Macros

GPU_DATA – (Make data accessible on GPU in specified region)

Macro Invocation

Uses FYPP call directive using #:call

#:call GPU_DATA(...)
   {code}
#:endcall GPU_DATA 

Parameters

name	data type	Default Value	description
code	code	Required	Region of code where defined data is accessible
copy	string list	None	Allocates and copies variable to GPU on entrance, then deallocated and copies to CPU on exit
copyin	string list	None	Allocates and copies data to GPU on entrance and then deallocated on exit
copyinReadOnly	string list	None	Allocates and copies a readonly variable to GPU and then deallocated on exit
copyout	string list	None	Allocates data on GPU on entrance and then deallocates and copies to CPU on exit
create	string list	None	Allocates data on GPU on entrance and then deallocates on exit
no_create	string list	None	Use data in CPU memory unless data is already in GPU memory (OpenACC only)
present	string list	None	Data that must be present in GPU memory. Increment counter on entrance, decrement on exit
deviceptr	string list	None	Pointer variables that are already allocated on GPU memory
attach	string list	None	Attaches device pointer to device targets on entrance, then detach on exit
default	string	None	Implicit assumptions compiler should make
extraAccArgs	string	None	String of any extra arguments added to the OpenACC directive
extraOmpArgs	string	None	String of any extra arguments added to the OpenMP directive

Parameter Restrictions

name	Restricted range
code	Do not assign it manually with key-value pairing

Example

#:call GPU_DATA(copy='[pixel_arr]', copyin='[starting_pixels, initial_index]',attach='[p_real, p_cmplx, p_fltr_cmplx]')
     {code}
     ...
#:endcall GPU_DATA
#:call GPU_DATA(create='[pixel_arr]', copyin='[initial_index]')
     {code}
     ...
#:endcall

GPU_ENTER_DATA – (Allocate/move data to GPU until matching GPU_EXIT_DATA or program termination)

Macro Invocation

Uses FYPP eval directive using $:

$:GPU_ENTER_DATA(...)

Parameter

name	data type	Default Value	description
copyin	string list	None	Allocates and copies data to GPU on entrance
copyinReadOnly	string list	None	Allocates and copies a readonly variable to GPU on entrance
create	string list	None	Allocates data on GPU on entrance
attach	string list	None	Attaches device pointer to device targets on entrance
extraAccArgs	string	None	String of any extra arguments added to the OpenACC directive
extraOmpArgs	string	None	String of any extra arguments added to the OpenMP directive

Example

$:GPU_ENTER_DATA(copyin='[pixels_arr]', copyinReadOnly='[starting_pixels, initial_index]')
$:GPU_ENTER_DATA(create='[bc_buffers(1:num_dims, -1:1)]', copyin='[initial_index]')

GPU_EXIT_DATA – (Deallocate/move data from GPU created by GPU_ENTER_DATA)

Macro Invocation

Uses FYPP eval directive using $:

$:GPU_EXIT_DATA(...)

Parameters

name	data type	Default Value	description
copyout	string list	None	Deallocates and copies data from GPU to CPU on exit
delete	string list	None	Deallocates data on GPU on exit
detach	string list	None	Detach device pointer from device targets on exit
extraAccArgs	string	None	String of any extra arguments added to the OpenACC directive
extraOmpArgs	string	None	String of any extra arguments added to the OpenMP directive

Example

$:GPU_EXIT_DATA(copyout='[pixels_arr]', delete='[starting_pixels, initial_index]')
$:GPU_EXIT_DATA(delete='[bc_buffers(1:num_dims, -1:1)]', copyout='[initial_index]')

GPU_DECLARE – (Allocate module variables on GPU or for implicit data region )

Macro Invocation

Uses FYPP eval directive using $:

$:GPU_DECLARE(...)

Parameters

name	data type	Default Value	description
copy	string list	None	Allocates and copies data to GPU on entrance, then deallocated and copies to CPU on exit
copyin	string list	None	Allocates and copies data to GPU on entrance and then deallocated on exit
copyinReadOnly	string list	None	Allocates and copies a readonly variable to GPU and then deallocated on exit
copyout	string list	None	Allocates data on GPU on entrance and then deallocates and copies to CPU on exit
create	string list	None	Allocates data on GPU on entrance and then deallocates on exit
present	string list	None	Data that must be present in GPU memory. Increment counter on entrance, decrement on exit
deviceptr	string list	None	Pointer variables that are already allocated on GPU memory
link	string list	None	Declare global link, and only allocate when variable used in data clause.
extraAccArgs	string	None	String of any extra arguments added to the OpenACC directive
extraOmpArgs	string	None	String of any extra arguments added to the OpenMP directive

Additional information

• An implicit data region is created at the start of each procedure and ends after the last executable statement in that procedure.
• Use only create, copyin, device_resident or link clauses for module variables
• GPU_DECLARE exit is the end of the implicit data region
• Link is useful for large global static data objects

Example

$:GPU_DECLARE(create='[x_cb,y_cb,z_cb,x_cc,y_cc,z_cc,dx,dy,dz,dt,m,n,p]')
$:GPU_DECLARE(create='[x_cb,y_cb,z_cb]', copyin='[x_cc,y_cc,z_cc]', link='[dx,dy,dz,dt,m,n,p]')

GPU_UPDATE – (Updates data from CPU to GPU or GPU to CPU)

Macro Invocation

Uses FYPP eval directive using $:

$:GPU_UPDATE(...)

Parameters

name	data type	Default Value	description
host	string list	None	Updates data from GPU to CPU
device	string list	None	Updates data from CPU to GPU
extraAccArgs	string	None	String of any extra arguments added to the OpenACC directive
extraOmpArgs	string	None	String of any extra arguments added to the OpenMP directive

Example

$:GPU_UPDATE(host='[arr1, arr2]')
$:GPU_UPDATE(host='[updated_gpu_val]', device='[updated_cpu_val]')

GPU_HOST_DATA – (Make GPU memory address available on CPU)

Macro Invocation

Uses FYPP call directive using #:call

#:call GPU_HOST_DATA(...)
   {code}
#:endcall GPU_HOST_DATA 

Parameters

name	data type	Default Value	description
code	code	Required	Region of code where GPU memory addresses is accessible
use_device_addr	string list	None	Use GPU memory address of variable instead of CPU memory address
use_device_ptr	string list	None	Use GPU pointer of pointers instead of CPU pointer
extraAccArgs	string	None	String of any extra arguments added to the OpenACC directive
extraOmpArgs	string	None	String of any extra arguments added to the OpenMP directive

Parameter Restrictions

name	Restricted range
code	Do not assign it manually with key-value pairing

Example

#:call GPU_HOST_DATA(use_device='[addr1, addr2]')
     {code}
     ...
#:endcall GPU_HOST_DATA
#:call GPU_HOST_DATA(use_device='[display_arr]')
     {code}
     ...
 #:endcall

---

Synchronization Macros

GPU_WAIT – (Makes CPU wait for async GPU activities)

Macro Invocation

Uses FYPP eval directive using $:

$:GPU_WAIT(...)

Parameters

name	data type	Default Value	description
extraAccArgs	string	None	String of any extra arguments added to the OpenACC directive
extraOmpArgs	string	None	String of any extra arguments added to the OpenMP directive

Example

$:GPU_WAIT()

GPU_ATOMIC – (Do an atomic operation on the GPU)

Macro Invocation

Uses FYPP eval directive using $:

$:GPU_ATOMIC(...)

Parameters

name	data type	Default Value	description
atomic	string	Required	Which atomic operation is performed
extraAccArgs	string	None	String of any extra arguments added to the OpenACC directive
extraOmpArgs	string	None	String of any extra arguments added to the OpenMP directive

Parameter Restrictions

name	Restricted range
atomic	'read', 'write', 'update', or 'capture'

Additional information

• read atomic is reading in a value
  • Ex: v=x
• write atomic is writing a value to a variable
  • Ex:x=square(tmp)
• update atomic is updating a variable in-place
  • Ex:x= x .and. 1

• Capture is a pair of read/write/update operations with one dependent on the other

  • Ex:

  x=x .and. 1
  v=x

Example

$:GPU_ATOMIC(atomic='update')
x = square(x)
$:GPU_ATOMIC(atomic='capture')
x = square(x)
v = x

---

Miscellaneous Macros

GPU_ROUTINE – (Compile a procedure for the GPU)

Macro Invocation

Uses FYPP eval directive using $:

$:GPU_ROUTINE(...)

Parameters

name	data type	Default Value	description
function_name	string	None	Name of subroutine/function
parallelism	string list	None	Parallelism granularity to use for this routine
nohost	boolean	False	Do not compile procedure code for CPU
cray_inline	boolean	False	Inline procedure on cray compiler
extraAccArgs	string	None	String of any extra arguments added to the OpenACC directive
extraOmpArgs	string	None	String of any extra arguments added to the OpenMP directive

Parameter Restrictions

name	Restricted range
parallelism	Valid elements: 'gang', 'worker', 'vector', 'seq'

Additional information

• Function name only needs to be given when cray_inline is True
• Future capability is to parse function header for function name
• Routine parallelism is most commonly '[seq]'

Example

$:GPU_ROUTINE(parallelism='[seq]')
$:GPU_ROUTINE(function_name='s_matmult', parallelism='[seq]', cray_inline=True)

GPU_CACHE – (Data to be cache in software-managed cache)

Macro Invocation

Uses FYPP eval directive using $:

$:GPU_CACHE(...)

Parameters

name	data type	Default Value	description
cache	string list	Required	Data that should to stored in cache
extraAccArgs	string	None	String of any extra arguments added to the OpenACC directive

NOTE Does not do anything for OpenMP currently

Example

$:GPU_CACHE(cache='[pixels_arr]')

---

Writing GPU_ROUTINE Device Helpers

GPU_ROUTINE marks a subroutine or function as callable from within a GPU kernel (an OpenACC routine or OpenMP declare target directive). The standard idiom for pure, sequential per-thread helpers is:

subroutine s_compute_mixture_coefficients(alpha_rho_K, alpha_K, rho_K, gamma_K, pi_inf_K, qv_K)
 
    $:gpu_routine(function_name='s_compute_mixture_coefficients', parallelism='[seq]', cray_inline=true)
 
    real(wp), dimension(num_fluids), intent(in) :: alpha_rho_K, alpha_K
    real(wp), intent(out)                       :: rho_K, gamma_K, pi_inf_K, qv_K
    ...
end subroutine s_compute_mixture_coefficients

When to use it. Extract a block into a GPU_ROUTINE helper when:

• It does roughly 50–100 or more FLOPs of work per call, enough to amortize the call overhead on device. Below that threshold, the helper is justified only when the same block appears verbatim in three or more callers (duplication removal).
• It is called from inside a GPU_PARALLEL_LOOP region — parallelism='[seq]' is correct for sequential per-thread work; the surrounding loop keeps the parallelism.

Key idioms.

• Always parallelism='[seq]' for these helpers. cray_inline=True is required when the helper is called from other modules — the Cray compiler does not inline cross-file routine calls without it, and performance collapses silently on that backend. Helpers called only from their own module (e.g., the shear/bulk stress tensor pair in m_riemann_state.fpp) do not need it.
• function_name= is required when cray_inline=True (the Cray inline directive needs the explicit name).
• The GPU_ROUTINE directive lives in the definition, not the call site. The caller needs no annotation beyond being inside a GPU_PARALLEL_LOOP.

Caller-loads, helper-computes. Callers do all coordinate-indexed array loads from the global state arrays (qL_rs_vf, qR_rs_vf, etc.) before the call; the helper receives only scalars or small arrays with explicit-shape dimensioning. This is required because the SF indexing lambda used in solver loops is defined locally inside each solver's #:for NORM_DIR block and cannot be referenced from a helper.

AMD case-opt compatibility. Under --case-optimization with the AMD backend, arrays that are sized by runtime parameters at compile time must be declared with an explicit constant bound. Use an explicit n argument (e.g., integer, intent(in) :: nf) and dimension helpers as dimension(nf) rather than dimension(num_fluids). See s_compute_interface_reynolds in src/simulation/m_riemann_state.fpp for the #:if not MFC_CASE_OPTIMIZATION and USING_AMD guard pattern: the guard sits on the dummy-argument declaration in the helper's definition, with matching guards on the callers' own local declarations so the actual and dummy bounds agree.

Declare scoping. The GPU_ROUTINE directive must appear in the source file that defines the routine. Helpers added to m_riemann_state.fpp are automatically in scope for every solver module that uses it — no additional declare-target annotations are needed at call sites.

Module boundaries and NVHPC inlining

Moving a device helper into another file can silently cost ~25% on NVHPC. NVHPC has no device LTO, so MFC's only cross-file inlining is the two-pass -Mextract=lib: / -Minline=lib: scheme in cmake/MFCTargets.cmake. That inliner refuses any device routine that has a subroutine call anywhere in its call tree, reporting:

subprogram not inlined -- missing prototype during crossing files: <name>

The refusal propagates: a caller of a refused routine is refused too. Measured on nvfortran 25.11 (A100), the following hold for a routine that must inline across files:

in the routine's body	inlines across files?
arithmetic, branches on module logicals, module array reads, early return	yes
a call to a scalar-returning function that is itself inlinable	yes (the callee need not be inlined)
a call to a subroutine	no
the routine returns a derived type or an array	no — never inlinable

There is no build-level escape: -Mextract always captures pre-inline source, so re-extracting in stages, compiling several files in one invocation, and levels:/maxsize:/name:/except: all fail, as does -Mipa (ignored in 25.x). Cray and AMD do their own whole-program IPA and are unaffected, and CPU builds do not care — so this shows up as an NVHPC-only benchmark regression while every other job stays green.

Symptom. Grind time regresses on NVHPC alone, with unchanged source semantics. Confirm by comparing per-routine stack frames: -Minfo=inline and -gpu=ptxinfo are already on, so Function properties for ... lines jumping from ~8 bytes to 200–350 bytes with matching spill stores/spill loads is the fingerprint. Registers per thread going down while the kernel gets slower is the same story seen from ncu.

Rule of thumb. Draw a module boundary where inlining already fails, not in the middle of a chain that currently inlines. Solver kernels never inlined s_compute_mixture_coefficients or s_compute_speed_of_sound even before m_eos existed, which makes that a free cut point; the phase chain those two call (s_phase_coefficients → s_eos_coefficients → s_reference_curve) must stay in the same file as them. This is why those four routines live in src/common/m_eos.fpp alongside the EOS operators even though mixture closure is not, strictly, an equation of state.

---

Debugging Tools and Tips for GPUs

Compiler agnostic tools

OpenMP tools

OMP_DISPLAY_ENV=true | false | verbose

• Prints out the internal control values and environment variables at the beginning of the program if true or verbose
• verbose will also print out vendor-specific internal control values and environment variables

OMP_TARGET_OFFLOAD = MANDATORY | DISABLED | DEFAULT

• Quick way to turn off off-load (DISABLED) or make it abort if a GPU isn't found (MANDATORY)
• Great first test: does the problem disappear when you drop back to the CPU?

OMP_THREAD_LIMIT=<positive_integer>

• Sets the maximum number of OpenMP threads to use in a contention group
• Might be useful in checking for issues with contention or race conditions

OMP_DISPLAY_AFFINITY=TRUE

• Will display affinity bindings for each OpenMP thread, containing hostname, process identifier, OS thread identifier, OpenMP thread identifier, and affinity binding.

Cray Compiler Tools

Cray General Options

CRAY_ACC_DEBUG: 0 (off), 1, 2, 3 (very noisy)

• Dumps a time-stamped log line (ACC: ...) for every allocation, data transfer, kernel launch, wait, etc. Great first stop when "nothing seems to run on the GPU".
• Outputs on STDERR by default. Can be changed by setting CRAY_ACC_DEBUG_FILE.
  • Recognizes stderr, stdout, and process.
  • process automatically generates a new file based on pid (each MPI process will have a different file)
• While this environment variable specifies ACC, it can be used for both OpenACC and OpenMP

CRAY_ACC_FORCE_EARLY_INIT=1

• Force full GPU initialization at program start so you can see start-up hangs immediately
• Default behavior without an environment variable is to defer initialization on first use
• Device initialization includes initializing the GPU vendor’s low-level device runtime library (e.g., libcuda for NVIDIA GPUs) and establishing all necessary software contexts for interacting with the device

Cray OpenACC Options

CRAY_ACC_PRESENT_DUMP_SAVE_NAMES=1

• Will cause acc_present_dump() to output variable names and file locations in addition to variable mappings
• Add acc_present_dump() around hotspots to help find problems with data movements
  • Helps more if adding CRAY_ACC_DEBUG environment variable

NVHPC Compiler Options

NVHPC General Options

STATIC_RANDOM_SEED=1

• Forces the seed returned by RANDOM_SEED to be constant, so it generates the same sequence of random numbers
• Useful for testing issues with randomized data

NVCOMPILER_TERM=option[,option]

• [no]debug: Enables/disables just-in-time debugging (debugging invoked on error)
• [no]trace: Enables/disables stack traceback on error

NVHPC OpenACC Options

NVCOMPILER_ACC_NOTIFY= <bitmask>

• Assign the environment variable to a bitmask to print out information to stderr for the following
  • kernel launches: 1
  • data transfers: 2
  • region entry/exit: 4
  • wait operation of synchronizations with the device: 8
  • device memory allocations and deallocations: 16
• 1 (kernels only) is the usual first step.3 (kernels + copies) is great for "why is it so slow?"

NVCOMPILER_ACC_TIME=1

• Lightweight profiler
• prints a tidy end-of-run table with per-region and per-kernel times and bytes moved
• Do not use with CUDA profiler at the same time

NVCOMPILER_ACC_DEBUG=1

• Spews everything the runtime sees: host/device addresses, mapping events, present-table look-ups, etc.
• Great for "partially present" or "pointer went missing" errors.
• Doc for NVCOMPILER_ACC_DEBUG
  • Ctrl+F for NVCOMPILER_ACC_DEBUG

NVHPC OpenMP Options

LIBOMPTARGET_PROFILE=run.json

• Emits a Chrome-trace (JSON) timeline you can open in chrome://tracing or Speedscope
• Great lightweight profiler when Nsight is overkill.
• Granularity in µs via LIBOMPTARGET_PROFILE_GRANULARITY (default 500).

LIBOMPTARGET_INFO=<bitmask>

• Prints out different types of runtime information
• Human-readable log of data-mapping inserts/updates, kernel launches, copies, waits.
• Perfect first stop for "why is nothing copied?"
• Flags
  • Print all data arguments upon entering an OpenMP device kernel: 0x01
  • Indicate when a mapped address already exists in the device mapping table: 0x02
  • Dump the contents of the device pointer map at kernel exit: 0x04
  • Indicate when an entry is changed in the device mapping table: 0x08
  • Print OpenMP kernel information from device plugins: 0x10
  • Indicate when data is copied to and from the device: 0x20

LIBOMPTARGET_DEBUG=1

• Developer-level trace (host-side)
• Much noisier than INFO
• Only works if the runtime was built with -DOMPTARGET_DEBUG.

LIBOMPTARGET_JIT_OPT_LEVEL=-O{0,1,2,3}

• This environment variable can be used to change the optimization pipeline used to optimize the embedded device code as part of the device JIT.
• The value corresponds to the -O{0,1,2,3} command line argument passed to clang.

LIBOMPTARGET_JIT_SKIP_OPT=1

• This environment variable can be used to skip the optimization pipeline during JIT compilation.
• If set, the image will only be passed through the backend.
• The backend is invoked with the LIBOMPTARGET_JIT_OPT_LEVEL flag.

AMD flang (amdflang) Known Issues

Whole-image device codegen instability (worked around in the build)

amdflang generates device code for the whole image at link time. Once the image carries enough OpenMP target regions, the device link's Attributor pass exceeds its AAPointerInfo access cap on a heavily shared object; pointer information goes pessimistic and OpenMPOpt's __kmpc_parallel cleanup then fails for the whole module. The visible effect: adding (or removing) ANY kernel anywhere silently regenerates UNTOUCHED kernels with far worse ISA — measured 2.4-4.5x slower, with register spills and an extra 512 B of LDS in every kernel. A wall-time A/B between two commits that differ in target-region count is confounded by this whole-image effect.

MFC's build raises the cap (-attributor-max-pi-accesses=16384, passed to the offload linker in cmake/MFCTargets.cmake), which restores full pointer precision for the whole image and makes kernel quality independent of unrelated edits. The cost is a longer device link. If a build's device link is unexpectedly slow, this flag is why — do not remove it; kernel performance becomes nondeterministic across commits without it.

The failure signature without the flag: after adding a kernel, unrelated kernels' resource usage shifts image-wide (uniform LDS increase, scratch/spill jumps visible in rocprofv3 dispatch records) and previously fast kernels slow several-fold.

Target regions inside Fortran BLOCK constructs are silently dropped

A GPU_PARALLEL_LOOP (OpenMP target region) written inside a Fortran block ... end block construct compiles cleanly, but amdflang omits it from the device image while the host still registers it. The first launch aborts with

hsa_executable_get_symbol_by_name(__omp_offloading_..._l<line>.kd):
HSA_STATUS_ERROR_INVALID_SYMBOL_NAME
omptarget error: Failed to load kernel ...

followed by a segmentation fault. Never place a GPU kernel inside a block construct; hoist it into its own (module) subroutine with the locals passed as arguments.

Silent-Failure Traps

Every entry here was measured. They share a failure mode: the build stays green and the answer is wrong, or one backend diverges from all the others.

• Do not wrap GPU_LOOP in GPU_PARALLEL for spatial loops. GPU_LOOP emits empty directives on Cray and AMD, so the loop runs serially with no diagnostic. Spatial loops always use GPU_PARALLEL_LOOP / END_GPU_PARALLEL_LOOP.
• An array whose bound is a device global (dimension(num_fluids), dimension(num_species)) may be passed to a device routine from a parallel-loop body, but not from inside another GPU_ROUTINE(parallelism='[seq]'). Cray OpenACC rejects the second form with ftn-7066 ... Global in accelerator routine without declare, and reports it at whatever line it gave up on: remove one trigger and the message walks forward to the next call, so the reported line is not the cause. Only the plain lanes fail, since --case-optimization turns those bounds into parameters — a green case-opt lane beside a failing plain one is the signature. Form such a call in the loop body and pass scalars deeper. Neither cray_inline, nor a num_fluids_max bound, nor dropping optional dummies avoids it; all three were tried.
• A device routine containing any GPU_LOOP must be called with scalars, never with an array element. On Cray OpenACC 19.0.0 through 21.0.2 at -O2 (-O0 and -O1 are correct, OpenMP offload is unaffected) the element is misaddressed: an intent(in) element reads as garbage and an intent(out) element is never written. Every routine level is affected, including a conforming loop vector inside routine vector. Either ingredient alone is fine, which is why a call like s_compute_pressure(q%sf(j,k,l), ...) into a loop-free helper works. Copy elements into locals before the call and receive into a local. Do not instead delete the seq directives: they are the idiom every device routine here uses. See #1815.
• Call m_thermochem species routines from the kernel, not from inside a GPU_ROUTINE. Calling get_species_* from within a device routine gives Cray OpenMP a runtime Memory access fault by GPU node-N on the first step while every other backend runs. The build is clean and only a case that reaches the path shows it. Evaluate them at the call site and pass the arrays in.
• nvfortran 23.11 and 24.1 segfault (fort2 TERMINATED by signal 11) on a caller that passes a parameter array from m_thermochem, such as molecular_weights, into a declare-target routine. Read such arrays directly in the kernel, or pass a plain local computed from them.
• The USING_AMD fypp guards are load-bearing, not a stale workaround. They swap a device-global array bound for a literal in src/common/include/shared_parallel_macros.fpp and its 86 use sites. Setting USING_AMD = False and rebuilding amdflang --gpu mp without case optimization compiles completely clean, then produces NaNs in CBC, the wave_speeds=2 Riemann path, immersed boundaries, surface tension, QBMM and viscous cases, and MHD HLLD, while both Lagrange bubble cases complete with out-of-tolerance answers. A compile-only check returns green, so any attempt to remove these must run the tests rather than just build.
• @:ACC_SETUP_VFs and @:ACC_SETUP_SFs compile only under Cray. Around MPI, use GPU_UPDATE(host=...) before a send and GPU_UPDATE(device=...) after a receive.

---

Compiler Documentation

• Cray & OpenMP Docs
• Cray & OpenACC Docs
• NVHPC & OpenACC Docs
• NVHPC & OpenMP Docs
• LLVM & OpenMP Docs
  • NVHPC is built on top of LLVM
• OpenMP Docs
• OpenACC Docs

Page last updated: 2026-02-04