This guide is intended for application programmers, scientists and engineers proficient in programming with the Fortran, C,
and/or C++ languages. The PGI tools are available on a variety of operating systems for the X86, AMD64, Intel 64, and OpenPOWER
hardware platforms. This guide assumes familiarity with basic operating system usage.

is used for filenames, directories, arguments, options, examples, and for language statements in the text, including assembly
language statements.

Bold

is used for commands.

[ item1 ]

in general, square brackets indicate optional items. In this case item1 is optional. In the context of p/t-sets, square brackets
are required to specify a p/t-set.

{ item2 | item 3 }

braces indicate that a selection is required. In this case, you must select either item2 or item3.

filename ...

ellipsis indicate a repetition. Zero or more of the preceding item may occur. In this example, multiple filenames are allowed.

FORTRAN

Fortran language statements are shown in the text of this guide using a reduced fixed point size.

C/C++

C/C++ language statements are shown in the test of this guide using a reduced fixed point size.

The PGI compilers and tools are supported on 64-bit variants of the Linux, Apple macOS, and Windows operating systems on a
variety of x86-compatible processors. There are a wide variety of releases and distributions of each of these types of operating
systems. The PGI compilers and PGI profiler are also supported on Linux for OpenPOWER.

Welcome to Release 2018 of PGI CUDA Fortran, a small set of extensions to Fortran that supports and is built upon the CUDA computing architecture.

Graphic processing units or GPUs have evolved into programmable, highly parallel computational units with very high memory
bandwidth, and tremendous potential for many applications. GPU designs are optimized for the computations found in graphics
rendering, but are general enough to be useful in many data-parallel, compute-intensive programs.

NVIDIA introduced CUDA™, a general purpose parallel programming architecture, with compilers and libraries to support the
programming of NVIDIA GPUs. CUDA comes with an extended C compiler, here called CUDA C, allowing direct programming of the
GPU from a high level language. The programming model supports four key abstractions: cooperating threads organized into thread
groups, shared memory and barrier synchronization within thread groups, and coordinated independent thread groups organized
into a grid. A CUDA programmer must partition the program into coarse grain blocks that can be executed in parallel. Each
block is partitioned into fine grain threads, which can cooperate using shared memory and barrier synchronization. A properly
designed CUDA program will run on any CUDA-enabled GPU, regardless of the number of available processor cores.

CUDA Fortran includes a Fortran 2003 compiler and tool chain for programming NVIDIA GPUs using Fortran. PGI 2018 includes support for CUDA Fortran on Linux, Apple macOS and Windows. CUDA Fortran is an analog to NVIDIA's CUDA C compiler.
Compared to the PGI Accelerator and OpenACC directives-based model and compilers, CUDA Fortran is a lower-level explicit programming
model with substantial runtime library components that give expert programmers direct control of all aspects of GPGPU programming.

The CUDA Fortran extensions described in this document allow the following operations in a Fortran program:

Declaring variables that are allocated in the GPU device memory

Allocating dynamic memory in the GPU device memory

Copying data from the host memory to the GPU memory, and back

Writing subroutines and functions to execute on the GPU

Invoking GPU subroutines from the host

Allocating pinned memory on the host

Using asynchronous transfers between the host and GPU

Using zero-copy and CUDA Unified Virtual Addressing features.

Accessing read-only data through texture memory caches.

Automatically generating GPU kernels using the kernel loop directive.

Launching GPU kernels from other GPU subroutines running on the device using CUDA 5.0 and above dynamic parallelism features.

Relocatable device code: Creating and linking device libraries such as the cublas; and calling functions defined in other modules and files.

In the CUDA Fortran host code on the left, device selection is explicit, performed by an API call on line 7. The provided cudafor module, used in line 2, contains interfaces to the full CUDA host runtime library, and in this case exposes the interface
to cudaSetDevice() and ensures it is called correctly. An array is allocated on the device at line 8. Line 9 of the host code initializes the
data on the host and the device, and, in line 10, a device kernel is launched. The interface to the device kernel is explicit,
in the Fortran sense, because the module containing the kernel is used in line 3. At line 11 of the host code, the results
from the kernel execution are moved back to a host array. Deallocation of the GPU array occurs on line 14.

Implicit Device Selection

Here is a CUDA Fortran example which is slightly more complicated than the preceding one.

In this case, the device selection is implicit, and defaults to NVIDIA device 0. The device array allocation in the host code at line 5 looks static, but actually occurs
at program init time. Larger array sizes are handled, both in the kernel launch at line 7 in the host code, and in the device
code at line 10. The device code contains examples of constant and shared data, which are described in Reference. There are actually two kernels launched from the host code: one explicitly provided and called from line 10, and a second,
generated using the CUDA Fortran kernel loop directive, starting at line 11. Finally, this example demonstrates the use of
the cublas module, used at line 2 in the host code, and called at line 12.

As these two examples demonstrate, all the steps listed at the beginning of this section for using a GPU are contained within
the host code. It is possible to program GPUs without writing any kernels and device code, through library calls and CUDA
Fortran kernel loop directives as shown, or by using higher-level directive-based models; however, programming in a lower-level
model like CUDA provides the programmer control over device resource utilization and kernel execution.

CUDA Fortran allows the definition of Fortran subroutines that execute in parallel on the GPU when called from the Fortran
program which has been invoked and is running on the host or, starting in CUDA 5.0, on the device. Such a subroutine is called
a device kernel or kernel.

A call to a kernel specifies how many parallel instances of the kernel must be executed; each instance will be executed by
a different CUDA thread. The CUDA threads are organized into thread blocks, and each thread has a global thread block index,
and a local thread index within its thread block.

A kernel is defined using the attributes(global) specifier on the subroutine statement; a kernel is called using special chevron syntax to specify the number of thread blocks
and threads within each thread block:

In this case, the call to the kernel ksaxpy specifies n/64 thread blocks, each with 64 threads. Each thread is assigned a thread block index accessed through the built-in blockidx variable, and a thread index accessed through threadidx. In this example, each thread performs one iteration of the common SAXPY loop operation.

Each thread is assigned a thread block index accessed through the built-in blockidx variable, and a thread index accessed through threadidx. The thread index may be a one-, two-, or three-dimensional index. In CUDA Fortran, the thread index for each dimension starts
at one.

Threads in the same thread block may cooperate by using shared memory, and by synchronizing at a barrier using the SYNCTHREADS() intrinsic. Each thread in the block waits at the call to SYNCTHREADS() until all threads have reached that call. The shared memory acts like a low-latency, high bandwidth software managed cache
memory. Currently, the maximum number of threads in a thread block is 1024.

A kernel may be invoked with many thread blocks, each with the same thread block size. The thread blocks are organized into
a one-, two-, or three-dimensional grid of blocks, so each thread has a thread index within the block, and a block index within the grid. When invoking a kernel,
the first argument in the chevron <<<>>> syntax is the grid size, and the second argument is the thread block size. Thread
blocks must be able to execute independently; two thread blocks may be executed in parallel or one after the other, by the
same core or by different cores.

The dim3 derived type, defined in the cudafor module, can be used to declare variables in host code which can conveniently hold the launch configuration values if they
are not scalars; for example:

CUDA Fortran programs have access to several memory spaces. On the host side, the host program can directly access data in
the host main memory. It can also directly copy data to and from the device global memory; such data copies require DMA access
to the device, so are slow relative to the host memory. The host can also set the values in the device constant memory, again
implemented using DMA access.

On the device side, data in global device memory can be read or written by all threads. Data in constant memory space is initialized
by the host program; all threads can read data in constant memory. Accesses to constant memory are typically faster than accesses
to global memory, but it is read-only to the threads and limited in size. Threads in the same thread block can access and
share data in shared memory; data in shared memory has a lifetime of the thread block. Each thread can also have private local
memory; data in thread local memory may be implemented as processor registers or may be allocated in the global device memory;
best performance will often be obtained when thread local data is limited to a small number of scalars that can be allocated
as processor registers.

Through use of the CUDA API as exposed by the cudafor module, access to CUDA features such as mapped memory, peer-to-peer memory access, and the unified virtual address space
are supported. Users should check the relevant CUDA documentation for compute capability restrictions for these features.
For an example of device array mapping, refer to Mapped Memory Example.

Starting with CUDA 6.0, managed or unified memory programming is available on certain platforms. For a complete description
of unified memory programming, see Appendix J. of the CUDA_C_Programming_Guide. Managed memory provides a common address space, and migrates data between the host and device as it is used by each set
of processors. On the host side, the data is resident in host main memory. On the device side, it is accessed as resident
in global device memory.

A subroutine or function in CUDA Fortran has an additional attribute, designating whether it is executed on the host or on
the device, and if the latter, whether it is a kernel, called from the host, or called from another device subprogram.

A subprogram declared with attributes(host), or with the host attribute by default, is called a host subprogram.

A subprogram declared with attributes(global) or attributes(device) is called a device subprogram.

A subroutine declared with attributes(global) is also called a kernel subroutine.

A subroutine declared with attributes(grid_global) is supported starting in CUDA 9.0 and on cc70 hardware or greater. Threads within the grid in these kernels are co-resident
on the same device and can be synchronized.

The host attribute, specified on the subroutine or function statement, declares that the subroutine or function is to be executed
on the host. Such a subprogram can only be called from another host subprogram. The default is attributes(host), if none of the host, global, or device attributes is specified.

The global attribute may only be specified on a subroutine statement; it declares that the subroutine is a kernel subroutine, to be
executed on the device, and may only be called using a kernel call containing the chevron syntax and runtime mapping parameters.

The device attribute, specified on the subroutine or function statement, declares that the subprogram is to be executed on the device;
such a routine must be called from a subprogram with the global or device attribute.

The grid_global attribute may only be specified on a subroutine statement; it declares that the subroutine is a kernel subroutine, to be
executed on the device, and may only be launched using a kernel call containing the chevron syntax and runtime mapping parameters.
The kernel is launched such that all threads within the grid group are guaranteed to be co-resident on the device. This allow
a grid synchronization operation on cc70 hardware and greater.

A device subprogram must not contain variables with the SAVE attribute, or with data initialization.

A kernel subroutine may not also have the device or host attribute.

Calls to a kernel subroutine must specify the execution configuration, as described in "Predefined Variables in Device Subprograms,"
on page 9. Such a call is asynchronous, that is, the calling routine making the call continues to execute before the device has completed its execution of the kernel
subroutine.

Device subprograms may not be contained in a host subroutine or function, and may not contain any subroutines or functions.

Variables in CUDA Fortran have a new attribute that declares in which memory the data is allocated. By default, variables
declared in modules or host subprograms are allocated in the host main memory. At most one of the device, managed, constant, shared, or pinned attributes may be specified for a variable.

A variable with the device attribute is called a device variable, and is allocated in the device global memory.

If declared in a module, the variable may be accessed by any subprogram in that module and by any subprogram that uses the
module.

If declared in a host subprogram, the variable may be accessed by that subprogram or subprograms contained in that subprogram.

A device array may be an explicit-shape array, an allocatable array, or an assumed-shape dummy array. An allocatable device
variable has a dynamic lifetime, from when it is allocated until it is deallocated. Other device variables have a lifetime
of the entire application.

Starting with CUDA 6.0, on certain platforms, a variable with the managed attribute is called a managed variable. Managed variables may be used in both host and device code. Variables with the managed attribute migrate between the host
and device, depending on where the accesses to the memory originate. Managed variables may be read and written by the host,
but there are access restrictions on the managed variables if kernels are active on the device. On the device, managed variables
have characteristics similar to device variables, but managed variables cannot be allocated from the device, as device variables
can be, starting in CUDA 5.0 in support of dynamic parallelism.

A variable with the constant attribute is called a device constant variable. Device constant variables are allocated in the device constant memory space. When declared in a module, the variable may
be accessed by any subprogram in that module and by any subprogram that uses the module. Device constant data may not be assigned
or modified in any device subprogram, but may be modified in host subprograms. Device constant variables may not be allocatable,
and have a lifetime of the entire application.

A variable with the shared attribute is called a device shared variable or a shared variable. A shared variable may only be declared in a device subprogram, and may only be accessed within that subprogram, or by other
device subprograms to which it is passed as an argument. A shared variable may not be data initialized. A shared variable
is allocated in the device shared memory for a thread block, and has a lifetime of the thread block. It can be read or written
by all threads in the block, though a write in one thread is only guaranteed to be visible to other threads after the next
call to the SYNCTHREADS() intrinsic.

A variable with the pinned attribute is called a pinned variable. A pinned variable must be an allocatable array. When a pinned variable is allocated, it will be allocated in host pagelocked
memory. The advantage of using pinned variables is that copies from page-locked memory to device memory are faster than copies
from normal paged host memory. Some operating systems or installations may restrict the use, availability, or size of page-locked
memory; if the allocation in page-locked memory fails, the variable will be allocated in the normal host paged memory and
required for asynchronous moves.

A variable with the texture attribute is called a texture variable. A texture variable must be an F90 pointer, and can be of type real or integer. Texture variables may be accessed only in
device subprograms, and can only be read, not written. The advantage of using texture variables is that the accesses to texture
data goes through a separate cache on the device, which may result in improved performance for many codes. Texture variables
are bound to underlying device arrays in host code using F90 pointer assignments.

Device subprograms have access to block and grid indices and dimensions through several built-in read-only variables. These
variables are of type dim3; the module cudafor defines the derived type dim3 as follows:

type(dim3)
integer(kind=4) :: x,y,z
end type

These predefined variables are not accessible in host subprograms.

The variable threadidx contains the thread index within its thread block; for one- or two-dimensional thread blocks, the threadidx%y and/or threadidx%z components have the value one.

The variable blockdim contains the dimensions of the thread block; blockdim has the same value for all thread blocks in the same grid.

The variable blockidx contains the block index within the grid; as with threadidx, for one-dimensional grids, blockidx%y and/or blockidx%z has the value one.

The variable griddim contains the dimensions of the grid.

The constant warpsize is declared to be type integer. Threads are executed in groups of 32, called warps; warpsize contains the number of threads in a warp, and is currently 32.

A call to a kernel subroutine must specify an execution configuration. The execution configuration defines the dimensionality
and extent of the grid and thread blocks that execute the subroutine. It may also specify a dynamic shared memory extent,
in bytes, and a stream identifier, to support concurrent stream execution on the device.

A kernel subroutine call looks like this:

call kernel<<<grid,block[,bytes][,streamid]>>>(arg1,arg2,...)

where

grid and block are either integer expressions (for one-dimensional grids and thread blocks), or are type(dim3), for one- or two-dimensional grids and thread blocks.

If grid is type(dim3), the value of each component must be equal to or greater than one, and the product is usually limited by the compute capability
of the device.

If block is type(dim3), the value of each component must be equal to or greater than one, and the product of the component values must be less than
or equal to 1024.

The value of bytes must be an integer; it specifies the number of bytes of shared memory to be allocated for each thread block, in addition
to the statically allocated shared memory. This memory is used for the assumed-size shared variables in the thread block;
refer to Shared data for more information. If the value of bytes is not specified, its value is treated as zero.

The value of streamid must be an integer greater than or equal to zero; it specifies the stream to which this call is associated. Nonzero stream
values can be created with a call to cudaStreamCreate. Starting in CUDA 7.0, the constant cudaStreamPerThread can be specified to use a unique default stream for each CPU thread.

When a host subprogram calls a kernel subroutine, the call actually returns to the host program before the kernel subroutine
begins execution. The call can be treated as a kernel launch operation, where the launch actually corresponds to placing the kernel on a queue for execution by the device. In this way,
the host can continue executing, including calling or queueing more kernels for execution on the device. By calling the runtime
routine cudaDeviceSynchronize, the host program can synchronize and wait for all previously launched or queued kernels.

Programmers must be careful when using concurrent host and device execution; in cases where the host program reads or modifies
device or constant data, the host program should synchronize with the device to avoid erroneous results.

Operations involving the device, including kernel execution and data copies to and from device memory, are implemented using
stream queues. An operation is placed at the end of the stream queue, and will only be initiated when all previous operations
on that queue have been completed.

An application can manage more concurrency by using multiple streams. Each user-created stream manages its own queue; operations
on different stream queues may execute out-of-order with respect to when they were placed on the queues, and may execute concurrently
with each other.

The default stream, used when no stream identifier is specified, is stream zero; stream zero is special in that operations
on the stream zero queue will begin only after all preceding operations on all queues are complete, and no subsequent operations
on any queue begin until the stream zero operation is complete.

CUDA Fortran allows automatic kernel generation and invocation from a region of host code containing one or more tightly nested
loops. Launch configuration and mapping of the loop iterations onto the hardware is controlled and specified as part of the
directive body using the familiar CUDA chevron syntax. As with any kernel, the launch is asynchronous. The program can use
cudaDeviceSynchronize() or CUDA Events to wait for the completion of the kernel.

The work in the loops specified by the directive is executed in parallel, across the thread blocks and grid; it is the programmer's
responsibility to ensure that parallel execution is legal and produces the correct answer. The one exception to this rule
is a scalar reduction operation, such as summing the values in a vector or matrix. For these operations, the compiler handles
the generation of the final reduction kernel, inserting synchronization into the kernel as appropriate.

The compiler maps the launch configuration specified by the grid and block values onto the outermost n loops, starting at loop n and working out. The grid and block values can be an integer scalar or a parenthesized list. Alternatively, using asterisks
tells the compiler to choose a thread block shape and/or compute the grid shape from the thread block shape and the loop limits.
Loops which are not mapped onto the grid and block values are run sequentially on each thread.

Kernel Loop Directive Example 1

In this example, the directive defines a two-dimensional thread block of size 32x4.

The body of the doubly-nested loop is turned into the kernel body:

ThreadIdx%x runs from 1 to 32 and is mapped onto the inner i loop.

ThreadIdx%y runs from 1 to 4 and is mapped onto the outer j loop.

The grid shape, specified as (*,*), is computed by the compiler and runtime by dividing the loop trip counts n and m by the thread block size, so all iterations are computed.

Kernel Loop Directive Example 2

!$cuf kernel do <<< *, 256 >>>
do j = 1, m
do i = 1, n
a(i,j) = b(i,j) + c(i,j)
end do
end do

Without an explicit n on the do, the schedule applies just to the outermost loop, that is, the default value is 1. In this case, only the outer j loop is run in parallel with a thread block size of 256. The inner i dimension is run sequentially on each thread.

You might consider if the code in Kernel Loop Directive Example 2 would perform better if the two loops were interchanged. Alternatively, you could specify a configuration like the following
in which the threads read and write the matrices in coalesced fashion.

Kernel Loop Directive Example 3

In Kernel Loop Directive Example 2, the 256 threads in each block each do one element of the matrix addition. Further expansion of the work along the i direction and all work across the j dimension is handled by the mapping onto the grid dimensions.

To "unroll" more work into each thread, specify non-asterisk values for the grid, as illustrated here:

Now the threads in a thread block handle all values in the i direction, in concert, incrementing by 256. One thread block is created for each j. Specifically, the j loop is mapped onto the grid x-dimension, because the compiler skips over the constant 1 in the i loop grid size. In CUDA built-in language, gridDim%x is equal to m.

If the directive specifies n dimensions, it must be followed by at least that many tightly-nested DO loops.

The tightly-nested DO loops must have invariant loop limits: the lower limit, upper limit, and increment must be invariant
with respect to any other loop in the kernel do.

There can be no GOTO or EXIT statements within or between any loops that have been mapped onto the grid and block configuration
values.

The body of the loops may contain assignment statements, IF statements, loops, and GOTO statements.

Only CUDA Fortran data types are allowed within the loops.

CUDA Fortran intrinsic functions are allowed, if they are allowed in device code, but the device-specific intrinsics such
as syncthreads, atomic functions, etc. are not.

Subroutine and function calls to attributes(device) subprograms are allowed if they are in the same module as the code containing
the directive.

Arrays used or assigned in the loop must have the device attribute.

Implicit loops and F90 array syntax are not allowed within the directive loops.

Scalars used or assigned in the loop must either have the device attribute, or the compiler will make a device copy of that
variable live for the duration of the loops, one for each thread. Except in the case of reductions; when a reduction has
a scalar target, the compiler generates a correct sequence of synchronized operations to produce one copy either in device
global memory or on the host.

Modern Fortran uses modules to package global data, definitions, derived types, and interface blocks. In CUDA Fortran these
modules can be used to easily communicate data and definitions between host and device code. This section includes a few examples
of using Fortran Modules.

In the host program, you use the top-level module, and get the definition of n and the interface to vadd. You can also rename the device arrays so they do not conflict with the host naming conventions:

Starting with CUDA 5.0, in addition to being able to access data declared in another module, you can also call device functions
which are contained in another module. In the following example, the file ffill.cuf contains a device function to fill an array:

Calling routines from other modules using relocatable device code.

To generate relocatable device code, compile this file with the -⁠Mcuda=rdc flag:

% pgf90 -Mcuda=rdc -c ffill.cuf

Now write another module and test program that calls the subroutine in this module. Since you are calling an attributes(device)
subroutine, you do not use the chevron syntax. For convenience, an overloaded Fortran sum function is included in the file
tfill.cuf which, in this case, takes 1-D integer device arrays.

Recently, PGI added support for F90 pointers that point to device data. Currently, this is limited to pointers that are declared
at module scope. The pointers can be accessed through module association, or can be passed in to global subroutines. The associated() function is also supported in device code. The following code shows many examples of using F90 pointers. These pointers can
also be used in CUF kernels.

In 2012, PGI added support for CUDA texture memory fetches through a special texture attribute ascribed to F90 pointers that
point to device data with the target attribute. In CUDA Fortran, textures are currently just for read-only data that travel
through the texture cache. Since there is separate hardware to support this cache, in many cases using the texture attribute
is a performance boost, especially in cases where the accesses are irregular and noncontiguous amongst threads. The following
simple example demonstrates this capability:

CUDA Fortran is supported by the PGI Fortran compilers when the filename
uses a CUDA Fortran extension. The .cuf extension
specifies that the file is a free-format CUDA Fortran program; the
.CUF extension may also be used, in which case the
program is processed by the preprocessor before being compiled. To
compile a fixed-format program, add the command line option
-⁠Mfixed. CUDA Fortran extensions can be enabled
in any Fortran source file by adding the -⁠Mcuda
command line option. It is important to remember that if you compile a
file with the -⁠Mcuda command line option, you
must also link the file with the -⁠Mcuda command
line option. If you compile with -⁠Mcuda, but do
not link with -⁠Mcuda, you will receive an
undefined reference to the symbol Mcuda_compiled.

To change the version of the CUDA Toolkit used from the default, specify
-⁠Mcuda=cudaX.Y; CUDA Toolkit version
X.Y must be installed.

Relocatable device code is generated by default. You can override this
option by specifying -⁠Mcuda=nordc.

If you are using many instances of the CUDA kernel loop directives, that
is, CUF kernels, you may want to add the -⁠Minfo
switch to verify that CUDA kernels are being generated where you expect
and whether you have followed the restrictions outlined in the preceding
sections.

Two GPU compiler code generators are supported, one based on CUDA C and
the other on LLVM. The LLVM backend is used by default; the CUDA C backend
can be enabled by the -⁠Mcuda=nollvm command line
option. Debug information generation for the GPU can only be generated by
the LLVM backend. To enable debugging on the host without using the
LLVM-based back-end, specify
-⁠g -⁠Mcuda=nodebug on the command line.

CUDA Fortran adds new attributes to subroutines and functions. This section describes how to specify the new attributes, their
meaning and restrictions.

A Subroutine may have the host, global, or device attribute, or may have both host and device attribute. A Function may have
the host or device attribute, or both. These attributes are specified using the attributes(attr) prefix on the Subroutine or Function statement; if there is no attributes prefix on the subprogram statement, then default
rules are used, as described in the following sections.

The host attributes prefix may be preceded or followed by any other allowable subroutine or function prefix specifiers (recursive,
pure, elemental, function return datatype). A subroutine or function with the host attribute is called a host subroutine or
function, or a host subprogram. A host subprogram is compiled for execution on the host processor. A subprogram with no attributes prefix has the host attribute
by default.

The global and grid_global attribute may be explicitly specified on the Subroutine statement as follows:

attributes(global) subroutine sub(...)

attributes(grid_global) subroutine subg(...)

Functions may not have a global attribute. A subroutine with either global attribute is called a kernel subroutine. A kernel subroutine may not be recursive, pure, or elemental, so no other subroutine prefixes are allowed. A kernel subroutine
is compiled as a kernel for execution on the device, to be called from a host routine using an execution configuration. A
kernel subroutine may not be contained in another subroutine or function, and may not contain any other subprogram. A grid_global
subroutine is supported starting with CUDA 9.0 on cc70 hardware or greater, and specifies that the kernel should be launched
in such a way that all threads in the grid can synchronize.

A subroutine or function with the device attribute may not be recursive, pure, or elemental, so no other subroutine or function
prefixes are allowed, except for the function return datatype. A subroutine or function with the device or kernel attribute
is called a device subprogram. A device subprogram is compiled for execution on the device. A subroutine or function with the device attribute must appear
within a Fortran module, and may only be called from device subprograms in the same module.

CUDA Fortran adds new attributes for variables and arrays. This section describes how to specify the new attributes and their
meaning and restrictions.

Variables declared in a host subprogram may have one of three new attributes: they may be declared to be in device global
memory, in managed memory, or in pinned memory.

Variables in modules may be declared to be in device global memory, in the managed memory space, or in constant memory space.
Additionally, the texture attribute can be added to read-only data declared in modules which enables reading the data through
the texture cache on the device.

Variables declared in a device program units may have one of three new attributes: they may be declared to be in device global
memory, in constant memory space, in the thread block shared memory, or without any additional attribute they will be allocated
in thread local memory. For performance and useability reasons, the value attribute can also be used on scalar dummy arguments
so they are passed by value, rather than the Fortran default to pass arguments by reference.

A variable or array with the device attribute is defined to reside in the device global memory. The device attribute can be
specified with the attributes statement, or as an attribute on the type declaration statement. The following example declares two arrays, a and b, to be device arrays of size 100.

Device variables and arrays may appear in modules, but may not be in a Common block or an Equivalence statement.

Members of a derived type may not have the device attribute unless they are allocatable.

Device variables and arrays may be passed as actual arguments to host and device subprograms; in that case, the subprogram
interface must be explicit (in the Fortran sense), and the matching dummy argument must also have the device attribute.

Device variables and arrays declared in a host subprogram cannot have the Save attribute unless they are allocatable.

In host subprograms, device data may only be used in the following manner:

In declaration statements

In Allocate and Deallocate statements

As an argument to the Allocated intrinsic function

As the source or destination in a data transfer assignment statement

As an actual argument to a kernel subroutine

As an actual argument to another host subprogram or runtime API call

As a dummy argument in a host subprogram

A device array may have the allocatable attribute, or may have adjustable extent.

A variable or array with the managed attribute is managed by the unified memory system and migrates between host main memory
and device global memory. The managed attribute can be specified with the attributes statement, or as an attribute on the type declaration statement. Managed arrays can be automatic or allocatable. The following
example declares two arrays, a and b, to be managed arrays of size 100, and allocates a third array, c with size 200.

Managed variables and arrays may appear in host subprograms and modules, but may not be in a Common block or an Equivalence
statement.

Managed variables and arrays declared in a host subprogram cannot have the Save attribute unless they are allocatable.

Derived types may have the managed attribute.

Members of a derived type may have the managed attribute.

Managed derived types may also contain allocatable device arrays.

Managed variables and arrays may be passed as actual arguments to other host subprograms; if the subprogram interface is overloaded,
the generic matching priority is match another managed dummy argument first, match a dummy with the device attribute next,
and match a dummy with no (or host) attribute last.

Passing a non-managed actual argument to a managed dummy argument will result in either a compilation error if the interface
is explicit, or unexpected behavior otherwise.

Managed variables and arrays may be passed as actual arguments to global subroutines just as device variables and arrays are.

By default, managed data is allocated with global scope, i.e. the flag passed to cudaMallocManaged is cudaMemAttachGlobal.

The scope of a managed variable can be changed with a call tocudaStreamAttachMemAsync.

Individual managed variables can be associated with a given stream by calling cudaforSetDefaultStream.

All subsequently allocated managed variables can also be associated with a given stream by calling cudaforSetDefaultStream.

Accessing managed data on the host while a running kernel is accessing managed data within the same scope on the device will
result in a segmentation fault.

These rules apply to managed data on the device:

The managed attribute may be used on dummy arguments.

Managed data is treated as if it were device data.

There is no support for allocating or deallocating managed data on the device.

Note:

Even if your application only uses a single GPU, if you are running on systems which have multiple GPUs that are not peer-to-peer
enabled, managed memory will be allocated as zero-copy memory and performance will suffer accordingly. A workaround is to
set the environment variable CUDA_VISIBLE_DEVICES so only one GPU is seen, or to force allocation on the GPU by setting CUDA_MANAGED_FORCE_DEVICE_ALLOC. The CUDA C Programming Guide has more details on this in Appendix J, section J.2.5.1.

An allocatable array with the pinned attribute will be allocated in special page-locked host memory, when such memory is available.
The advantage of using pinned memory is that transfers between the device and pinned memory are faster and can be asynchronous.
An array with the pinned attribute may be declared in a module or in a host subprogram. The pinned attribute can be specified
with the attributes statement, or as an attribute on the type declaration statement. The following example declares two arrays, p and q, to be pinned allocatable arrays.

Pinned arrays may be passed as arguments to host subprograms regardless of whether the interface is explicit, or whether the
dummy argument has the pinned and allocatable attributes. Where the array is deallocated, the declaration for the array must
still have the pinned attribute, or the deallocation may fail.

A variable or array with the constant attribute is defined to reside in the device constant memory space. The constant attribute can be specified with the attributes statement, or as an attribute on the type declaration statement. The following example declares two arrays, c and d, to be constant arrays of size 100.

real :: c(100)
attributes(constant) :: c
real, constant :: d(100)

These rules apply to constant data:

Constant variables and arrays can appear in modules, but may not be in a Common block or an Equivalence statement. Constant variables appearing in modules may be accessed via the use statement in both host and device subprograms.

Constant data may not have the Pointer, Target, or Allocatable attributes.

Members of a derived type may not have the constant attribute.

Arrays with the constant attribute must have fixed size.

Constant variables and arrays may be passed as actual arguments to host and device subprograms, as long as the subprogram
interface is explicit, and the matching dummy argument also has the constant attribute. Constant variables cannot be passed as actual arguments between a host subprogram and a device global subprogram.

Within device subprograms, variables and arrays with the constant attribute may not be assigned or modified.

Within host subprograms, variables and arrays with the constant attribute may be read and written.

In host subprograms, data with the constant attribute may only be used in the following manner:

A variable or array with the shared attribute is defined to reside in the shared memory space of a thread block. A shared
variable or array may only be declared and used inside a device subprogram. The shared attribute can be specified with the
attributes statement, or as an attribute on the type declaration statement. The following example declares two arrays, s and t, to be shared arrays of size 100.

real :: c(100)
attributes(shared) :: c
real, shared :: d(100)

These rules apply to shared data:

Shared data may not have the Pointer, Target, or Allocatable attributes.

Shared variables may not be in a Common block or Equivalence statement.

Members of a derived type may not have the shared attribute.

Shared variables and arrays may be passed as actual arguments to from a device subprogram to another device subprogram, as
long as the interface is explicit and the matching dummy argument has the shared attribute.

Shared arrays that are not dummy arguments may be declared as assumed-size arrays; that is, the last dimension of a shared
array may have an asterisk as its upper bound:

real, shared :: x(*)

Such an array has special significance. Its size is determined at run time by the call to the kernel. When the kernel is called,
the value of the bytes argument in the execution configuration is used to specify the number of bytes of shared memory that is dynamically allocated
for each thread block. This memory is used for the assumed-size shared memory arrays in that thread block; if there is more
than one assumed-size shared memory array, they are all implicitly equivalenced, starting at the same shared memory address.
Programmers must take this into account when coding.

Shared arrays may be declared as Fortran automatic arrays. For automatic arrays, the bounds are declared as an expression
containing constants, parameters, blockdim variables, and integer arguments passed in by value. The allocation of automatic
arrays also comes from the dynamic area specified via the chevron launch configuration. If more than one automatic array
is declared, the compiler and runtime manage the offsets into the dynamic area. Programmers must provide a sufficient number
of bytes in the chevron launch configuration shared memory value to cover all automatic arrays declared in the global subroutine.

If a shared array is not a dummy argument and not assumed-size or automatic, it must be fixed size. In this case, the allocation
for the shared array does not come from the dynamically allocated shared memory area specified in the launch configuration,
but rather it is declared statically within the function. If the global routine uses only fixed size shared arrays, or none
at all, no shared memory amount needs to be specified at the launch.

Read-only real and integer device data can be accessed in device subprograms through the texture memory by assigning an F90
pointer variable to the underlying device array. To use texture memory in this manner, follow these steps:

Add a declaration to a module declaration section that is used in both the host and device code:

real, texture, pointer :: t(:)

In your host code, add the target attribute to the device data that you wish to access via texture memory:

Change: real, device :: a(n)

To: real, target, device :: a(n)

The target attribute is standard F90/F2003 syntax to denote an array or other data structure that may be "pointed to" by another
entity.

Tie the texture declaration to the device array by using the F90 pointer assignment operator in your host code. A simple expression
like the following one performs all the underlying CUDA texture binding operations.

t => a

The CUDA Fortran device code that can refer to t through use or host association can now access the elements of t without any change in syntax.

In the following example, accesses of t, targeting a, go through the texture cache.

In device subprograms, following the rules of Fortran, dummy arguments are passed by default by reference. This means the
actual argument must be stored in device global memory, and the address of the argument is passed to the subprogram. Scalar
arguments can be passed by value, as is done in C, by adding the value attribute to the variable declaration.

In this case, the value of n can be passed from the host without needing to reside in device memory. The variable arrays corresponding
to the dummy arguments a and b must be set up before the call to reside on the device.

Device arrays can have the allocatable attribute. These arrays are dynamically allocated in host subprograms using the Allocate
statement, and dynamically deallocated using the Deallocate statement. If a device array declared in a host subprogram does
not have the Save attribute, it will be automatically deallocated when the subprogram returns.

Scalar variables can be allocated on the device using the Fortran 2003 allocatable scalar feature. To use these, declare and
initialize the scalar on the host as:

integer, allocatable, device :: ndev
allocate(ndev)
ndev = 100

The language also supports the ability to create the equivalent of automatic and local device arrays without using the allocate
statement. These arrays will also have a lifetime of the subprogram as is usual with the Fortran language:

For programmers comfortable with the CUDA C programming environment, Fortran interfaces to the CUDA memory management runtime
routines are provided. These functions return memory which will bypass certain Fortran allocatable properties such as automatic
deallocation, and thus the arrays are treated more like C malloc’ed areas. Mixing standard Fortran allocate/deallocate with
the runtime Malloc/Free for a given array is not supported.

The cudaMalloc function can be used to allocate single-dimensional arrays of the supported intrinsic data-types, and cudaFree
can be used to free it:

Allocatable arrays with the pinned attribute are dynamically allocated using the Allocate statement. The compiler will generate
code to allocate the array in host page-locked memory, if available. If no such memory space is available, or if it is exhausted,
the compiler allocates the array in normal paged host memory. Otherwise, pinned allocatable arrays work and act like any other
allocatable array on the host.

CUDA Fortran supports the ability to create the equivalent of automatic and local managed arrays without using the allocate
statement. These arrays will also have a lifetime of the subprogram as is usual with the Fortran language:

You can copy variables and arrays from the host memory to the device memory by using simple assignment statements in host
subprograms. By default, using assignment statements to read or write device, managed, or constant data implicitly uses CUDA
stream zero. This means such data copies are synchronous, and the data copy waits until all previous kernels and data copies
complete. Alternatively, you can use the cudaforSetDefaultStream call to associate one or more device and managed variables to a particular stream. After this call has occurred, assignment
statements on those variables will run asynchronously on the specified stream.

Specific information on assignment statements:

An assignment statement where the left hand side is a device variable or device array or array section, and the right hand
side is a host variable or host array or array section, copies data from the host memory to the device global memory.

An assignment statement where the left hand side is a host variable or host array or array section, and the right hand side
is a device variable or device array or array section, copies data from the device global memory to the host memory.

An assignment statement with a device variable or device array or array section on both sides of the assignment statement
copies data between two device variables or arrays.

Similarly, you can use simple assignment statements to copy or assign variables or arrays with the constant attribute.

Specific information on assignment statements and managed data:

An assignment statement where the left hand side is a managed variable or managed array, and the right hand side is a conforming
scalar constant, host variable, host array or array section, copies data from the host memory to the device global memory
using cudaMemcpy, memset, or a similar operation.

An assignment statement where the left hand side is a managed array section and the right hand side is any host variable copies
data using generated host code.

An assignment statement where the left hand side is a managed variable, managed array or array section, and the right hand
side is a device variable or device array or array section, copies data from the device global memory to the host memory using
cudaMemcpy or a similar operation.

An assignment statement where the right hand side is a managed variable or managed array, and the left hand side is a host
variable, host array or array section, copies data from the device global memory to the host memory using cudaMemcpy or a
similar operation.

An assignment statement where the right hand side is a managed array section and the left hand side is any host or managed
variable copies data using generated host code.

An assignment statement where the right hand side is a managed variable, managed array or array section, and the left hand
side is a device variable or device array or array section, copies data using cudaMemcpy and accesses the data from the device.

More information on Memcpy and Memset behavior with managed memory can be found in Appendix J. of the CUDA_C_Programming_Guide.

Some limited data transfer can be enclosed within expressions. In general, the rule of thumb is all arithmetic or operations
must occur on the host, which normally only allows one device array to appear on the right-hand-side of an expression. Temporary
arrays are generated to accommodate the host copies of device data as needed. For instance, if a, b, and c are conforming host arrays, and adev, bdev, and cdev are conforming device arrays, the following expressions are legal:

a = adev

adev = a

b = a + adev

c = x * adev + b

The following expressions are not legal as they either promote a false impression of where the actual computation occurs,
or would be more efficient written in another way, or both:

c = adev + bdev

adev = adev + a

b = sqrt(adev)

Elemental transfers are supported by the language but perform poorly. Array slices are also supported, and their performance
is dependent on the size of the slice, the amount of contiguous data in the slices, and the implementation.

For programmers comfortable with the CUDA C programming environment, Fortran interfaces to the CUDA memory management runtime
routines are provided. These functions can transfer data either from the host to device, device to host, or from one device
array to another.

The cudaMemcpy function can be used to copy data between the host and the GPU:

For those familiar with the CUDA C routines, the kind parameter to the Memcpy routines is optional in Fortran because the
attributes of the arrays are explicitly declared. Counts expressed in arguments to the Fortran runtime routines are expressed
in terms of data type elements, not bytes.

For a complete list of memory management runtime routines, refer to Memory Management.

A call to a kernel subroutine must give the execution configuration for the call. The execution configuration gives the size
and shape of the grid and thread blocks that execute the function as well as the amount of shared memory to use for assumed-size
shared memory arrays and the associated stream.

The execution configuration is specified after the subroutine name in the call statement; it has the form:

<<< grid, block, bytes, stream >>>

grid is an integer, or of type(dim3). If it is type(dim3), the value of grid%z must be one. The product grid%x*grid%y gives the number of thread blocks to launch. If grid is an integer, it is converted to dim3(grid,1,1). bl

block is an integer, or of type(dim3). If it is type(dim3), the number of threads per thread block is block%x*block%y*block%z, which must be less than the maximum supported by the device. If block is an integer, it is converted to dim3(block,1,1).

bytes is optional; if present, it must be a scalar integer, and specifies the number of bytes of shared memory to be allocated
for each thread block to use for assumed-size shared memory arrays. For more information, refer to Shared Data. If not specified, the value zero is used.

stream is optional; if present, it must be an integer, and have a value of zero, or a value returned by a call to cudaStreamCreate.
See Section 4.5 on page 41. It specifies the stream to which this call is enqueued. Beginning with PGI 15.4, the stream constant
value cudaStreamPerThread may be specified. This will use a unique stream for each CPU thread. Also, starting in PGI 15.4, the default stream value,
normally zero, can be changed by a call to cudaforSetDefaultStream().

For instance, a kernel subroutine

attributes(global) subroutine sub( a )

can be called like:

call sub <<< DG, DB, bytes >>> ( A )

The function call fails if the grid or block arguments are greater than the maximum sizes allowed, or if bytes is greater than the shared memory available. Shared memory may also be consumed by fixed-sized shared memory declarations
in the kernel and for other dedicated uses, such as function arguments and execution configuration arguments.

Variables and arrays with the device, constant, or shared attributes, or declared in device subprograms, are limited to the
types described in this section. They may have any of the intrinsic datatypes in the following table.

Table 2. Device Code Intrinsic Datatypes

Type

Type Kind

integer

1,2,4(default),8

logical

1,2,4(default),8

real

4(default),8

double precision

equivalent to real(kind=8)

complex

4(default),8

character(len=1)

1 (default)

Additionally, they may be of derived type, where the members of the derived type have one of the allowed intrinsic datatypes,
or another allowed derived type.

The system module cudafor includes definitions of the derived type dim3, defined as

The variable threadidx contains the thread index within its thread block; for one- or two-dimensional thread blocks, the threadidx%y and/or threadidx%z components have the value one.

The variable blockdim contains the dimensions of the thread block; blockdim has the same value for all threads in the same grid; for one- or two-dimensional thread blocks, the blockdim%y and/or blockdim%z components have the value one.

The variable blockidx contains the block index within the grid; as with threadidx, for one-dimensional grids, blockidx%y has the value one. The value of blockidx%z is always one. The value of blockidx is the same for all threads in the same thread block.

The variable griddim contains the dimensions of the grid; the value of griddim%z is always one. The value of griddim is the same for all threads in the same grid; the value of griddim%z is always one; the value of griddim%y is one for one-dimensional grids.

The variables threadidx, blockdim, blockidx, and griddim are available only in device subprograms.

The constant warpsize contains the number of threads in a warp. It is currently defined to be 32.

SYNCTHREADS

The syncthreads intrinsic subroutine acts as a barrier synchronization for all threads in a single thread block; it has no arguments:

void syncthreads()

Sometimes threads within a block access the same addresses in shared or global memory, thus creating potential read-after-write,
write-after-read, or write-after-write hazards for some of these memory accesses. To avoid these potential issues, use syncthreads()to specify synchronization points in the kernel. This intrinsic acts as a barrier at which all threads in the block must wait
before any thread is allowed to proceed. Threads within a block cooperate and share data by synchronizing their execution
to coordinate memory accesses.

Each thread in a thread block pauses at the syncthreads call until all threads have reached that call. If any thread in a thread block issues a call to syncthreads, all threads must also reach and execute the same call statement, or the kernel fails to complete correctly.

SYNCTHREADS_AND

integer syncthreads_and(int_value)

syncthreads_and. like syncthreads, acts as a barrier at which all threads in the block must wait before any thread is allowed to proceed. In addition, syncthreads_and evaluates the integer argument int_value for all threads of the block and returns non-zero if and only if int_value evaluates to non-zero for all of them.

SYNCTHREADS_COUNT

integer syncthreads_count(int_value)

syncthreads_count, like syncthreads, acts as a barrier at which all threads in the block must wait before any thread is allowed to proceed. In addition, syncthreads_count evaluates the integer argument int_value for all threads of the block and returns the number of threads for which int_value evaluates to non-zero.

SYNCTHREADS_OR

integer syncthreads_or(int_value)

syncthreads_or. like syncthreads, acts as a barrier at which all threads in the block must wait before any thread is allowed to proceed. In addition, syncthreads_or evaluates the integer argument int_value for all threads of the block and returns non-zero if and only if int_value evaluates to non-zero for any of them.

Memory Fences

In general, when a thread issues a series of writes to memory in a particular order, other threads may see the effects of
these memory writes in a different order. You can use threadfence(), threadfence_block(), and threadfence_system() to create a memory fence to enforce ordering.

For example, suppose you use a kernel to compute the sum of an array of N numbers in one call. Each block first sums a subset
of the array and stores the result in global memory. When all blocks are done, the last block done reads each of these partial
sums from global memory and sums them to obtain the final result. To determine which block is finished last, each block atomically
increments a counter to signal that it is done with computing and storing its partial sum. If no fence is placed between storing
the partial sum and incrementing the counter, the counter might increment before the partial sum is stored.

THREADFENCE

void threadfence()

threadfence acts as a memory fence, creating a wait. Typically, when a thread issues a series of writes to memory in a particular order,
other threads may see the effects of these memory writes in a different order. threadfence() is one method to enforce a specific order. All global and shared memory accesses made by the calling thread prior to threadfence() are visible to:

All threads in the thread block for shared memory accesses

All threads in the device for global memory accesses

THREADFENCE_BLOCK

void threadfence_block()

threadfence_block acts as a memory fence, creating a wait until all global and shared memory accesses made by the calling thread prior to threadfence_block() are visible to all threads in the thread block for all accesses.

THREADFENCE_SYSTEM

void threadfence_system()

threadfence_system acts as a memory fence, creating a wait until all global and shared memory accesses made by the calling thread prior to threadfence_system() are visible to:

All threads in the thread block for shared memory accesses

All threads in the device for global memory accesses

Host threads for page-locked host memory accesses

threadfence_system() is only supported by devices of compute capability 2.0 or higher.

Warp-vote operations are only supported by devices with compute capability 1.2 and higher. Each of these functions has a single
argument.

ALLTHREADS

The allthreads function is a warp-vote operation with a single scalar logical argument:

if( allthreads(a(i)<0.0) ) allneg = .true.

The function allthreads evaluates its argument for all threads in the current warp. The value of the function is .true. only if the value of the argument is .true. for all threads in the warp.

ANYTHREAD

The anythread function is a warp-vote operation with a single scalar logical argument:

if( anythread(a(i)<0.0) ) allneg = .true.

The function anythread evaluates its argument for all threads in the current warp. The value of the function is .false. only if the value of the argument is .false. for all threads in the warp.

BALLOT

The ballot function is a warp-vote operation with a single integer argument:

unsigned integer ballot(int_value)

The function ballot evaluates the argument int_value for all threads of the warp and returns an integer whose Nth bit is set if and only if int_value evaluates to non-zero for the Nth thread of the warp.

The atomic functions read and write the value of their first operand,
which must be a variable or array element in shared memory (with the
shared attribute) or in device global memory (with the device
attribute). Atomic functions are only supported by devices with compute
capability 1.1 and higher. Compute capability 1.2 or higher is required if
the first argument has the shared attribute. Certain real(4) and real(8)
atomic functions may require compute capability 2.0 and higher.

The atomic functions return correct values even if multiple threads in the
same or different thread blocks try to read and update the same location
without any synchronization.

Arithmetic and Bitwise Atomic Functions

These atomic functions read and return the value of the first
argument. They also combine that value with the value of the second
argument, depending on the function, and store the combined value back
to the first argument location. For atomicadd, atomicsub, atomicmax,
atomicmin, and atomicexch, the data types may be integer(4), integer(8),
real(4), or real(8). For atomicand, atomicor, and atomicxor, only
integer(4) arguments are supported.

Note:
The return value for each of these functions is the first argument,
mem.

These functions are:

Table 9. Arithmetic and Bitwise Atomic Functions

Function

Additional Atomic Update

atomicadd( mem, value )

mem = mem + value

atomicsub( mem, value )

mem = mem − value

atomicmax( mem, value )

mem = max(mem,value)

atomicmin( mem, value )

mem = min(mem,value)

atomicand( mem, value )

mem = iand(mem,value)

atomicor( mem, value )

mem = ior(mem,value)

atomicxor( mem, value )

mem = ieor(mem,value)

atomicexch( mem, value )

mem = value

Counting Atomic Functions

These atomic functions read and return the value of the first
argument. They also compare the first argument with the second argument,
and stores a new value back to the first argument location, depending on
the result of the comparison. These functions are intended to implement
circular counters, counting up to or down from a maximum value specified
in the second argument. Both arguments must be of type integer(kind=4).

Note:
The return value for each of these functions is the first argument,
mem.

These functions are:

Table 10. Counting Atomic Functions

Function

Additional Atomic Update

atomicinc( mem, imax )

if (mem<imax) then
mem = mem+1
else
mem = 0
endif

atomicdec( mem, imax )

if (mem<imax .and. mem>0) then
mem = mem-1
else
mem = imax
endif

Compare and Swap Atomic Function

This atomic function reads and returns the value of the first
argument. It also compares the first argument with the second argument,
and atomically stores a new value back to the first argument location if
the first and second argument are equal. All three arguments must be of
the same type, either integer(kind=4), integer(kind=8), real(kind=4), or
real(kind=8).

Starting in PGI 15.1, improvements have been made in the support of list-directed PRINT or WRITE statements to the default
output unit (PRINT * or WRITE(*,*)). Before PGI 15.1, the output for PRINT or WRITE statements would likely be interleaved
between different threads for each item on the PRINT or WRITE statement. That is, if a device kernel contains a PRINT statement,
such as this one:

print *, 'index = ', i, j

the user would observe the character string 'index = ', one after the other from all threads, then the second item, the value
of i, then the third item, j, and finally the end-of-line.

Unlike the usual C printf implementation, which prints out a whole line for each thread, there was no indication of which
thread prints out which item, and in which order.

Now, starting with PGI 15.1, the PGI Fortran accelerator runtime, shared between CUDA Fortran and OpenACC for NVIDIA targets,
buffers up the output and prints an entire line's worth in one operation. In PGI 15.1 through 15.4, the integer and character
types were supported. Beginning in PGI 15.5, logical, real and complex types are also supported.

The underlying CUDA printf implementation limits the number of print statements in a kernel launch to 4096. Users should
take this limit into account when making use of this feature.

Starting in PGI 15.5, print and write statements in device code are also supported when used with the LLVM code generator
-ta=llvm, or -Mcuda=llvm, and in combination with the -mp compiler option.

By adding the compiler option -Mcuda=charstring, some limited support for character strings, character substrings, character variables, and string assignment is also now
available in CUDA Fortran device code. Here is a short example:

__shfl()

__shfl() returns the value of var held by the thread whose ID is given by srcLane. If the srcLane is outside the range of 1:width, then the thread's own value of var is returned. The width argument is optional in all shuffle functions and has a default value of 32, the current warp size.

__shfl_up()

__shfl_up() calculates a source lane ID by subtracting delta from the caller's thread ID. The value of var held by the resulting thread ID is returned; in effect, var is shifted up the warp by delta lanes.

The source lane index will not wrap around the value of width, so the lower delta lanes are unchanged.

__shfl_down()

__shfl_down() calculates a source lane ID by adding delta to the caller's thread ID. The value of var held by the resulting thread ID is returned: this has the effect of shifting var down the warp by delta lanes. The ID number of the source lane will not wrap around the value of width, so the upper delta lanes remain unchanged.

__shfl_xor()

__shfl_xor() uses ID-1 to calculate the source lane ID by performing a bitwise XOR of the caller's lane ID with the laneMask. The value of var held by the resulting lane ID is returned. If the resulting lane ID falls outside the range permitted by width, the thread's own value of var is returned. This mode implements a butterfly addressing pattern such as is used in tree reduction and broadcast.

Beginning in PGI 15.1, the sum, maxval, and minval host intrinsics are overloaded to accept device or managed arrays when the cudafor module is used. The tuned kernels written
to implement these operations use the shuffle functions, so compute capability 3.0 or higher is required to use these reductions.
If the mask optional argument is used, the mask argument must be either a device logical array, or an expression containing
managed operands and constants, i.e. the mask must be computable on the host but readable on the device.

Here is a complete example which performs the sum and maxval reductions on the GPU:

By default, intrinsic reductions that are supported on the device will be executed on the device for (large enough) managed
arrays. There may be occasions where one would like to perform reductions on managed data on the host. This can be accomplished
using the rename feature of the "use" statement, for example:

By default, these reductions will run on a nonzero stream. If a unique per-thread default stream was set via a call to cudaforSetDefaultStream, the reduction initialization will pick that up. Users can further control which stream reductions run on, or force stream
0, by using the cudaforReductionSetStream() call. The current reduction stream can be queried using the cudaforReductionGetStream() call.

SUM() on complex managed or device data is not currently supported. This support will be added in a future PGI release

CUDA Fortran Modules are available to help programmers access features of the CUDA runtime environment, which might otherwise
not be accessible from Fortran without significant effort from the programmer. These modules might be either device modules
or host modules.

PGI provides a device module by default which allows access and interfaces to many of the CUDA device built-in routines.

To access this module explicitly, do one of the following:

Add this line to your Fortran program:

use cudadevice

Add this line to your C program:

#include <cudadevice.h>

You can use these routines in CUDA Fortran global and device subprograms, in CUF kernels, and in PGI Accelerator compute regions
in Fortran as well as in C. Further, the PGI compilers come with implementations of these routines for host code, though these
implementations are not specifically optimized for the host.

On NVIDIA GPUs which support CUDA Compute Capability 7.0 and above, PGI provides a device module which provides interfaces
to cooperative group functionality which is provided by NVIDIA starting in CUDA 9.0.

To access this module, add this line to your Fortran subprogram:

use cooperative_groups

Here is a simple example of using the cooperative_groups device module which enables a cooperative grid kernel:

There is currently limited functionality for cooperative groups of size less than or equal to a thread block. More functionality
will be added in an upcoming release. Currently, the following types are defined within the module: grid_group, thread_group, and coalesced_group. Each type has two public members, the size and rank. The syncthreads subroutine is overloaded in the cooperative_groups
module to take the type as an argument, to appropriately synchronize the threads in that group. Minimal code sequences supported
are:

PGI provides a module which defines interfaces to the CUBLAS Library from PGI CUDA Fortran. These interfaces are made accessible
by placing the following statement in the CUDA Fortran host-code program unit.

use cublas

The interfaces are currently in three forms:

Overloaded traditional BLAS interfaces which take device arrays as arguments rather than host arrays, i.e.

New CUBLAS 4.0+ interfaces with access to all features of the new library.

These interfaces are all in the form of function calls, take a handle as the first argument, and pass many scalar arguments
and results by reference, i.e.

istat = cublasSaxpy_v2(h, n, a, x, incx, y, incy)

In the case of saxpy, users now have the option of having "a" reside either on the host or device. Functions which traditionally return a scalar, such as sdot() and isamax(), now take an extra argument for returning the result. Functions which traditionally take a character*1 argument, such as 't' or 'n' to control transposing, now take an integer value defined in the cublas module.

To support the third form, a derived type named cublasHandle is defined in the cublas module. You can define a variable of this type using

type(cublasHandle) :: h

Initialize it by passing it to the cublasCreate function.

When using CUBLAS 4.0 and higher, the cublas module properly generates handles for the first two forms from serial and OpenMP
parallel regions.

Intermixing the three forms is permitted. To access the handles used internally in the cublas module use:

h = cublasGetHandle()

The following form "istat = cublasGetHandle(h)" is also supported.

istat = cublasGetHandle(h)

Assignment and tests for equality and inequality are supported for the cublasHandle type.

Refer to Cublas Module Example for an example that demonstrates the use of the cublas module, the cublasHandle type, and the three forms of calls.

PGI provides another module which defines interfaces to the CUFFT Library from PGI CUDA Fortran. These interfaces are made
accessible by placing the following statement in the CUDA Fortran host-code program unit.

Beginning in PGI 15.1, the distribution also contains a module which defines interfaces to the CUSPARSE Library from PGI CUDA
Fortran. The CUSPARSE library has had fairly substantial changes made to it over the last several CUDA releases, so the source
for the module is also included in our package, and a different version of the pre-compiled module is put into a version-specific
location which is searched according to the CUDA version specified on the compilation command line. These interfaces are
made explicit by placing the following statement in the CUDA Fortran host-code program unit.

use cusparse

In addition to the function interfaces, there are several important derived types and constants which are defined in the cusparse
module. Here is an example of their use:

The system module cudafor defines the interfaces to the Runtime API routines.

Most of the runtime API routines are integer functions that return an error code; they return a value of zero if the call
was successful, and a nonzero value if there was an error. To interpret the error codes, refer to Error Handling.

cudaDeviceGetCacheConfig returns the preferred cache configuration for the current device. Current possible cache configurations are defined to be
cudaFuncCachePreferNone, cudaFuncCachePreferShared, and cudaFuncCachePreferL1.

cudaDeviceGetCacheConfig is available in device code starting in CUDA 5.0.

cudaDeviceGetSharedMemConfig returns the current size of the shared memory banks on the current device. This routine is for use with devices with configurable
shared memory banks, and is supported starting with CUDA 4.2. Current possible shared memory configurations are defined to
be cudaSharedMemBankSizeDefault, cudaSharedMemBankSizeFourByte, and cudaSharedMemBankSizeEightByte.

cudaDeviceSetCacheConfig sets the current device preferred cache configuration. Current possible cache configurations are defined to be cudaFuncCachePreferNone, cudaFuncCachePreferShared, and cudaFuncCachePreferL1.

cudaDeviceSetSharedMemConfig sets the size of the shared memory banks on the current device. This routine is for use with devices with configurable shared
memory banks, and is supported starting with CUDA 4.2. Current possible shared memory configurations are defined to be cudaSharedMemBankSizeDefault, cudaSharedMemBankSizeFourByte, and cudaSharedMemBankSizeEightByte.

Sometimes threads within a block access the same addresses in shared or global memory, thus creating potential read-after-write,
write-after-read, or write-after-write hazards for some of these memory accesses. To avoid these potential issues, use the
functions in this section for thread management. These functions have been deprecated beginning in CUDA 4.0.

cudaforGetDefaultStream returns the default stream which has been associated with a thread, managed variable, or device variable via a call to cudaforSetDefaultStream. devptr may be any managed or device scalar or array of a supported type specified in Device Code Intrinsic Datatypes. The devptr argument is optional; if it is not specified, the function returns the stream tied to the thread, or zero (the default stream).

cudaGetStreamDefault was available starting in CUDA 6.0. In PGI 15.4, it was renamed to cudaforGetDefaultStream since this function is not available in CUDA C/C++.

Streams values returned from cudaforGetDefaultStream can be used as the argument to other CUDA libraries, such as the routines cublasSetStream(), cufftSetStream(), and cusparseSetStream().

cudaforSetDefaultStream sets the default stream for all subsequent high-level CUDA Fortran operations on managed or device data initiated by that
CPU thread. Starting in PGI 15.4, the specific operations affected with managed data are allocatation via the Fortran allocate
statement, assignment (both memset and memcpy types), CUF Kernel and global kernel launches, and sum(), maxval(), and minval()
reduction operations. devptr may be any managed or device scalar or array of a supported type specified in Device Code Intrinsic Datatypes. The devptr argument is optional; if it is not specified, the function ties the specified stream to all subsequent, allowable, high-level
operations executing on that thread.

cudaSetStreamDefault was available starting in CUDA 6.0. In PGI 15.4, it was renamed to cudaforSetDefaultStream since this function is not available in CUDA C/C++.

cudaStreamAttachMemAsync initiates a stream operation to attach the managed allocation starting at address devptr to the specified stream. devptr may be any managed scalar or array of a supported type specified in Device Code Intrinsic Datatypes. The argument len is optional, but currently must be zero. The flags argument must be cudaMemAttachGlobal, cudMemAttachHost, or cudMemAttachSingle.

cudaStreamCreateWithPriority creates an asynchronous stream and assigns its identifier to its first argument. Valid values for flags are cudaStreamDefault or cudaStreamNonBlocking. Lower values for priority represent higher priorities. Work in a higher priority stream may preempt work already executing
in a low priority stream.

cudaStreamQuery tests whether all operations enqueued to the selected stream are complete; it returns zero (success) if all operations are complete, and the value cudaErrorNotReady if not. It may also return another error condition if some asynchronous operations failed.

cudaStreamSynchronize blocks execution of the host subprogram until all preceding kernels and operations associated with the given stream are complete. It may return error codes from previous, asynchronous operations.

cudaEventElapsedTime computes the elapsed time between two events (in milliseconds). It returns cudaErrorInvalidValue if either event has not yet been recorded. This function is only valid with events recorded on stream zero.

cudaEventQuery tests whether an event has been recorded. It returns success (zero) if the event has been recorded, and cudaErrorNotReady if it has not. It returns cudaErrorInvalidValue if cudaEventRecord has not been called for this event.

cudaEventRecord issues an operation to the given stream to record an event. The event is recorded after all preceding operations in the stream are complete. If stream is zero, the event is recorded after all preceding operations in all streams are complete.

cudaFuncSetCacheConfig sets the preferred cache configuration for the function named by the func argument, which must be a global function. Current
possible cache configurations are defined to be cudaFuncCachePreferNone, cudaFuncCachePreferShared, and cudaFuncCachePreferL1.

cudaFuncSetSharedMemConfig sets the size of the shared memory banks for the function named by the func argument, which must be a global function. This routine is for use with devices with configurable shared memory banks, and
is supported starting with CUDA 4.2. Current possible shared memory configurations are defined to be cudaSharedMemBankSizeDefault, cudaSharedMemBankSizeFourByte, and cudaSharedMemBankSizeEightByte

Many of the memory management routines can take device arrays as arguments. Some can also take C types, provided through the
Fortran 2003 iso_c_binding module, as arguments to simplify interfacing to existing CUDA C code.

CUDA Fortran has extended the F2003 derived type TYPE(C_PTR) by providing a C device pointer, defined in the cudafor module, as TYPE(C_DEVPTR). Consistent use of TYPE(C_PTR) and TYPE(C_DEVPTR), as well as consistency checks between Fortran device arrays and host arrays, should be of benefit.

Currently, it is possible to construct a Fortran device array out of a TYPE(C_DEVPTR) by using an extension of the iso_c_binding subroutine c_f_pointer. Under CUDA Fortran, c_f_pointer will take a TYPE(C_DEVPTR) as the first argument, an allocatable device array as the second argument, a shape as the third argument, and in effect transfer
the allocation to the Fortran array. Similarly, there is also a function C_DEVLOC() defined which will create a TYPE(C_DEVPTR) that holds the C address of the Fortran device array argument. Both of these features are subject to change when, in the
future, proper Fortran pointers for device data are supported.

cudaHostAlloc allocates pinned memory on the host. It returns in hostptr the address of the page-locked allocation, or returns an error if the memory is unavailable. Size is in bytes. The flags argument enables different options to be specified that affect the allocation. The normal iso_c_binding subroutine c_f_pointer
can be used to move the type(c_ptr) to a Fortran pointer.

cudaHostGetDevicePointer returns a pointer to a device memory address corresponding to the pinned memory on the host. hostptr is a pinned memory buffer that was allocated via cudaHostAlloc(). It returns in devptr an address that can be passed to, and read and written by, a kernel which runs on the device. The flags argument is provided for future releases. The normal iso_c_binding subroutine c_f_pointer can be used to move the type(c_devptr)to
a device array.

cudaMalloc allocates data on the device. devptr may be any allocatable, one-dimensional device array of a supported type specified in Device Code Intrinsic Datatypes. The count is in terms of elements. Or, devptr may be of TYPE(C_DEVPTR), in which case the count is in bytes.

cudaMallocManaged allocates data that will be managed by the unified memory system. devptr may be any allocatable, one-dimensional managed array of a supported type specified in Device Code Intrinsic Datatypes. The count is in terms of elements. Or, devptr may be of TYPE(C_DEVPTR), in which case the count is in bytes. The flags argument must be either cudaMemAttachGlobal or cudaMemAttachHost.

cudaMallocPitch allocates data on the device. devptr may be any allocatable, two-dimensional device array of a supported type specified in Device Code Intrinsic Datatypes. The width is in terms of number of elements. The height is an integer.

cudaMallocPitch may pad the data, and the padded width is returned in the variable pitch. Pitch is an integer of kind=cuda_count_kind. devptr may also be of TYPE(C_DEVPTR), in which case the integer values are expressed in bytes.

cudaMalloc3D allocates data on the device. pitchptr is a derived type defined in the cudafor module. cext is also a derived type which holds the extents of the allocated array. Alternatively, pitchptr may be any allocatable, three-dimensional device array of a supported type specified in Datatypes Allowed.

cudaMemAdvise lends advice to the Unified Memory subsystem about the expected usage pattern for the specified memory range. devptr may be any managed memory scalar or array, of a supported type specified in Device Code Intrinsic Datatypes. The count is in terms of elements. Alternatively, devptr may be of TYPE(C_DEVPTR), in which case the count is in terms of bytes.

Current possible values for advice, defined in the cudafor module, are cudaMemAdviseSetReadMostly, cudaMemAdviseUnsetReadMostly, cudaMemAdviseSetPreferredLocation, cudaMemAdviseUnsetPreferredLocation, cudaMemAdviseSetAccessedBy, and cudaMemAdviseUnsetAccessedBy

The device argument specifies the destination device. Passing in cudaCpuDeviceId for the device, which is defined as a parameter in the cudafor module, will set advice for the CPU.

cudaMemcpy copies data from one location to another. dst and src may be any device or host, scalar or array, of a supported type specified in Device Code Intrinsic Datatypes. The count is in terms of elements. kdir may be optional; for more information, refer to Data Transfer Using Runtime Routines. If kdir is specified, it must be one of the defined enums cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, or cudaMemcpyDeviceToDevice. Alternatively, dst and src may be of TYPE(C_DEVPTR) or TYPE(C_PTR), in which case the count is in term of bytes.

cudaMemcpyAsync copies data from one location to another. dst and src may be any device or host, scalar or array, of a supported type specified in Device Code Intrinsic Datatypes. The count is in terms of elements. kdir may be optional; for more information, refer to Data Transfer Using Runtime Routines. If kdir is specified, it must be one of the defined enums cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, or cudaMemcpyDeviceToDevice. Alternatively, dst and src may be of TYPE(C_DEVPTR) or TYPE(C_PTR), in which case the count is in term of bytes.

This function operates on page-locked host memory only. The copy can be associated with a stream by passing a non-zero stream
argument; otherwise the stream argument is optional and defaults to zero.

cudaMemcpyFromSymbol copies data from a device area in global or constant memory space referenced by a symbol to a destination on the host. dst may be any host scalar or array of a supported type specified in Datatypes Allowed. The count is in terms of elements.

cudaMemcpyFromSymbolASYNC copies data from a device area in global or constant memory space referenced by a symbol to a destination on the host. dst may be any host scalar or array of a supported type specified in Datatypes Allowed. The count is in terms of elements.

cudaMemcpyFromSymbolASYNCis asynchronous with respect to the host, This function operates on page-locked host memory only. The copy can be associated
with a stream by passing a non-zero stream argument.

cudaMemcpyPeer copies data from one device to another. dst and src may be any device scalar or array, of a supported type specified in Device Code Intrinsic Datatypes. The count is in terms of elements. Alternatively, dst and src may be of TYPE(C_DEVPTR), in which case the count is in term of bytes.

cudaMemcpyPeerAsync copies data from one device to another. dst and src may be any device scalar or array, of a supported type specified in Device Code Intrinsic Datatypes. The count is in terms of elements. Alternatively, dst and src may be of TYPE(C_DEVPTR), in which case the count is in term of bytes. The copy can be associated with a stream by passing a non-zero stream argument.

cudaMemcpyToSymbol copies data from the source to a device area in global or constant memory space referenced by a symbol. src may be any host scalar or array of a supported type as specified in Device Code Intrinsic Datatypes. The count is in terms of elements.

cudaMemcpyToSymbolAsync copies data from the source to a device area in global or constant memory space referenced by a symbol. src may be any host scalar or array of a supported type specified in Datatypes Allowed. The count is in terms of elements.

This function operates on page-locked host memory only. The copy can be associated with a stream by passing a non-zero stream
argument.

cudaMemcpy2D copies data from one location to another. dst and src may be any device or host array, of a supported type specified in Device Code Intrinsic Datatypes. The width and height are in terms of elements. kdir may be optional; for more information, refer to Data Transfer Using Runtime Routines. If kdir is specified, it must be one of the defined enums cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, or cudaMemcpyDeviceToDevice. Alternatively, dst and src may be of TYPE(C_DEVPTR) or TYPE(C_PTR), in which case the width and height are in term of bytes.

cudaMemcpy2D copies data from one location to another. dst and src may be any device or host array, of a supported type specified in Device Code Intrinsic Datatypes. The width and height are in terms of elements. kdir may be optional; for more information, refer to Data Transfer Using Runtime Routines. If kdir is specified, it must be one of the defined enums cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, or cudaMemcpyDeviceToDevice. Alternatively, dst and src may be of TYPE(C_DEVPTR) or TYPE(C_PTR), in which case the width and height are in term of bytes.

This function operates on page-locked host memory only. The copy can be associated with a stream by passing a non-zero stream argument, otherwise the stream argument is optional and defaults to zero.

cudaMemPrefetchAsync prefetches memory to the specified destination device. devptr may be any managed memory scalar or array, of a supported type specified in Device Code Intrinsic Datatypes. The count is in terms of elements. Alternatively, devptr may be of TYPE(C_DEVPTR), in which case the count is in terms of bytes.

The device argument specifies the destination device.
The stream argument specifies which stream to enqueue the prefetch operation on.

Passing in cudaCpuDeviceId for the device, which is defined as a parameter in the cudafor module, will prefetch the data to CPU memory.

cudaMemset sets a location or array to the specified value. devptr may be any device scalar or array of a supported type specified in Device Code Intrinsic Datatypes. The value must match in type and kind. The count is in terms of elements. Or, devptr may be of TYPE(C_DEVPTR), in which case the count is in term of bytes, and the lowest byte of value is used.

cudaMemsetAsync sets a location or array to the specified value. devptr may be any device scalar or array of a supported type specified in Device Code Intrinsic Datatypes. The value must match in type and kind. The count is in terms of elements. Or, devptr may be of TYPE(C_DEVPTR), in which case the count is in term of bytes, and the lowest byte of value is used. The memory set operation is associated with the stream specified.

cudaMemset2D sets an array to the specified value. devptr may be any device array of a supported type specified in Device Code Intrinsic Datatypes. The value must match in type and kind. The pitch, width, and height are in terms of elements. Or, devptr may be of TYPE(C_DEVPTR), in which case the pitch, width, and height are in terms of bytes, and the lowest byte of value is used.

cudaPointerGetAttributes returns the attributes of a device or host pointer in the attributes type. ptr may be any host or device scalar or array of a supported type specified in Datatypes Allowed. It may also be of type C_PTR or C_DEVPTR.

This example shows a program to compute the product C of two matrices A and B, as follows:

Each thread block computes one 16x16 submatrix of C;

Each thread within the block computes one element of the submatrix.

The submatrix size is chosen so the number of threads in a block is a multiple of the warp size (32) and is less than the
maximum number of threads per thread block (512).

Each element of the result is the product of one row of A by one column of B. The program computes the products by accumulating
submatrix products; it reads a block submatrix of A and a block submatrix of B, accumulates the submatrix product, then moves
to the next submatrix of A rowwise and of B columnwise. The program caches the submatrices of A and B in the fast shared memory.

For simplicity, the program assumes the matrix sizes are a multiple of 16, and has not been highly optimized for execution
time.

This source code module mmul_mod has two subroutines. The host subroutine mmul is a wrapper for the kernel routine mmul_kernel.

MMUL

This host subroutine has two input arrays, A and B, and one output array, C, passed as assumed-shape arrays. The routine performs the following operations:

It determines the size of the matrices in N, M, and L.

It allocates device memory arrays Adev, Bdev, and Cdev.

It copies the arrays A and B to Adev and Bdev using array assignments.

It fills dimGrid and dimBlock to hold the grid and thread block sizes.

It calls mmul_kernel to compute Cdev on the device.

It copies Cdev back from device memory to C.

It frees the device memory arrays.

Because the data copy operations are synchronous, no extra synchronization is needed between the copy operations and the kernel
launch.

MMUL_KERNEL

This kernel subroutine has two device memory input arrays, A and B, one device memory output array, C, and three scalars giving the array sizes. The thread executing this routine is one of 16x16 threads cooperating in a thread
block. This routine computes the dot product of A(i,:)*B(:,j) for a particular value of i and j, depending on the block and thread index.

It performs the following operations:

It determines the thread indices for this thread.

It determines the i and j indices, for which element of C(i,j) it is computing.

It initializes a scalar in which it will accumulate the dot product.

It steps through the arrays A and B in blocks of size 16.

For each block, it does the following steps:

It loads one element of the submatrices of A and B into shared memory.

It synchronizes to make sure both submatrices are loaded by all threads in the block.

It accumulates the dot product of its row and column of the submatrices.

It synchronizes again to make sure all threads are done reading the submatrices before starting the next block.

This example demonstrates the use of CUDA API supported in the cudafor module for mapping page-locked host memory into the address space of the device. It makes use of the iso_c_binding c_ptr type and the cudafor c_devptr types to interface to the C routines, then the Fortran c_f_pointer call to map the types to Fortran arrays.

This example demonstrates the use of the cublas module, the cublasHandle type, the three forms of cublas calls, and the use of mapped pinned memory, all within the framework of an multi-threaded
OpenMP program.

This example demonstrates the use of CUDA managed memory in an OpenMP program. In the main program, one stream is created
for each OpenMP thread. A call to cudaforSetDefaultStream is made to set that as the default stream for all subsequent high-level
language constructs. The default stream is used explicitly in the launch configuration of the CUF kernel, and also as the
thread's input argument for synchronization. Once the cudaStreamSynchronize has occurred, this thread can safely access the
managed data on the host, in this case in the any() function, even while other threads may be in the middle of their kernel
launch.

The PGI User
Forum is monitored
by members of the PGI engineering and support teams as well as other
PGI customers. The forums contain answers to many commonly asked questions.
Log in
to the PGI website to access the forums.

Notice

ALL NVIDIA DESIGN SPECIFICATIONS, REFERENCE BOARDS, FILES, DRAWINGS, DIAGNOSTICS, LISTS, AND OTHER DOCUMENTS (TOGETHER AND
SEPARATELY, "MATERIALS") ARE BEING PROVIDED "AS IS." NVIDIA MAKES NO WARRANTIES, EXPRESSED, IMPLIED, STATUTORY, OR OTHERWISE
WITH RESPECT TO THE MATERIALS, AND EXPRESSLY DISCLAIMS ALL IMPLIED WARRANTIES OF NONINFRINGEMENT, MERCHANTABILITY, AND FITNESS
FOR A PARTICULAR PURPOSE.

Information furnished is believed to be accurate and reliable. However, NVIDIA Corporation assumes no responsibility for the
consequences of use of such information or for any infringement of patents or other rights of third parties that may result
from its use. No license is granted by implication of otherwise under any patent rights of NVIDIA Corporation. Specifications
mentioned in this publication are subject to change without notice. This publication supersedes and replaces all other information
previously supplied. NVIDIA Corporation products are not authorized as critical components in life support devices or systems
without express written approval of NVIDIA Corporation.

Trademarks

NVIDIA, the NVIDIA logo, Cluster Development Kit, PGC++, PGCC, PGDBG, PGF77, PGF90, PGF95, PGFORTRAN, PGHPF, PGI, PGI Accelerator,
PGI CDK, PGI Server, PGI Unified Binary, PGI Visual Fortran, PGI Workstation, PGPROF, PGROUP, PVF, and The Portland Group
are trademarks and/or registered trademarks of NVIDIA Corporation in the U.S. and other countries. Other company and product
names may be trademarks of the respective companies with which they are associated.