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 following restrictions apply to subprograms.

• 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.

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.

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.

The general form of the kernel directive is:

    !$cuf kernel do[(n)] <<< grid, block  [optional stream] >>>

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.

There are two ways to specify the optional stream argument:

    !$cuf kernel do[(n)] <<< grid, block, 0, streamid >>>

Or

    !$cuf kernel do[(n)] <<< grid, block, stream=streamid >>>

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

    !$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

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

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

The following restrictions apply to CUF kernel directives:

• 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.

The simplest directive form for performing a dot product on two device arrays takes advantage of the properties for scalar use outlined previously.

    rsum = 0.0
    !$cuf kernel do <<< *, * >>>
    do i = 1, n
        rsum = rsum + x(i)* y(i)
    end do

For reductions, the compiler recognizes the use of the scalar and generates just one final result.

This CUF kernel can be followed by another CUF kernel in the same subprogram:

    !$cuf kernel do <<< *, * >>>
    do i = 1, n
        rsum= x(i) * y(i)
        z(i) = rsum
    end do

In this CUF kernel, the compiler recognizes rsum as a scalar temporary which should be allocated locally on every thread. However, use of rsum on the host following this loop is undefined.

in the following example, a set of modules are defined in one file which are accessed by another module.

module moda
  real, device, allocatable :: a(:)
end module

module modb
  real, device, allocatable :: b(:)
end module

module modc
  use moda
  use modb
  integer, parameter :: n = 100
  real, device, allocatable :: c(:)
  contains
    subroutine vadd()
    !$cuf kernel do <<<*,*>>>
    do i = 1, n
      c(i) = a(i) + b(i)
    end do
    end subroutine
end module

program t
use modc, a_d => a, b_d => b, c_d => c
real a,b,c(n)
allocate(a_d(n),b_d(n),c_d(n))
a_d = 1.0
b_d = 2.0
call vadd()
c = c_d
print *,all(c.eq.3.0)
end

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:

module ffill
  contains
    attributes(device) subroutine fill(a)
    integer, device :: a(*)
    i = (blockidx%x-1)*blockdim%x + threadidx%x
    a(i) = i
    end subroutine
end module

% pgf90 -Mcuda=rdc -c ffill.cuf

module testfill 
    use ffill
    contains
    attributes(global) subroutine Kernel(arr)
        integer, device :: arr(*)
        call fill(arr)
    end subroutine Kernel

    integer function sum(arr)
        integer, device :: arr(:)
        sum = 0
        !$cuf kernel do <<<*,*>>>
        do i = 1, size(arr)
          sum = sum + arr(i)
        end do
    end function sum
end module testfill

program tfill
use testfill
integer, device :: iarr(100)
iarr = 0
call Kernel<<<1,100>>>(iarr)
print *,sum(iarr)==100*101/2
end program tfill

% pgf90 -Mcuda=rdc tfill.cuf ffill.o

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.

module devptr
! currently, pointer declarations must be in a module
  real, device, pointer, dimension(:) :: mod_dev_ptr
  real, device, pointer, dimension(:) :: arg_dev_ptr
  real, device, target,  dimension(4) :: mod_dev_arr
  real, device, dimension(4) :: mod_res_arr
contains
  attributes(global) subroutine test(arg_ptr)
    real, device, pointer, dimension(:) :: arg_ptr
    ! copy 4 elements from one of two spots
    if (associated(arg_ptr)) then
      mod_res_arr = arg_ptr
    else
      mod_res_arr = mod_dev_ptr
    end if
  end subroutine test
end module devptr

program test
use devptr
real, device, target, dimension(4) :: a_dev
real  result(20)

a_dev = (/ 1.0, 2.0, 3.0, 4.0 /)

! Pointer assignment to device array declared on host,
! passed as argument.  First four result elements.
arg_dev_ptr => a_dev
call test<<<1,1>>>(arg_dev_ptr)
result(1:4) = mod_res_arr

!$cuf kernel do <<<*,*>>>
do i = 1, 4
  mod_dev_arr(i) = arg_dev_ptr(i) + 4.0
  a_dev(i)       = arg_dev_ptr(i) + 8.0
end do

! Pointer assignment to module array, argument nullified
! Second four result elements
mod_dev_ptr => mod_dev_arr
arg_dev_ptr => null()
call test<<<1,1>>>(arg_dev_ptr)
result(5:8) = mod_res_arr

! Pointer assignment to updated device array, now associated
! Third four result elements
arg_dev_ptr => a_dev
call test<<<1,1>>>(arg_dev_ptr)
result(9:12) = mod_res_arr

!$cuf kernel do <<<*,*>>>
do i = 1, 4
  mod_dev_arr(i) = 25.0 - mod_dev_ptr(i)
  a_dev(i)       = 25.0 - arg_dev_ptr(i)
end do

! Non-contiguous pointer assignment to updated device array
! Fourth four element elements
arg_dev_ptr => a_dev(4:1:-1)
call test<<<1,1>>>(arg_dev_ptr)
result(13:16) = mod_res_arr

! Non-contiguous pointer assignment to updated module array
! Last four elements of the result
nullify(arg_dev_ptr)
mod_dev_ptr => mod_dev_arr(4:1:-1)
call test<<<1,1>>>(arg_dev_ptr)
result(17:20) = mod_res_arr

print *,all(result==(/(real(i),i=1,20)/))
end

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:

module memtests
  real(8), texture, pointer :: t(:)  ! declare the texture
  contains
    attributes(device) integer function bitrev8(i)
    integer ix1, ix2, ix
    ix = i
    ix1 = ishft(iand(ix,z'0aa'),-1)
    ix2 = ishft(iand(ix,z'055'), 1)
    ix = ior(ix1,ix2)
    ix1 = ishft(iand(ix,z'0cc'),-2)
    ix2 = ishft(iand(ix,z'033'), 2)
    ix = ior(ix1,ix2)
    ix1 = ishft(ix,-4)
    ix2 = ishft(ix, 4)
    bitrev8 = iand(ior(ix1,ix2),z'0ff')
    end function bitrev8
    
    attributes(global) subroutine without( a, b )
    real(8), device :: a(*), b(*)
    i = blockDim%x*(blockIdx%x-1) + threadIdx%x
    j = bitrev8(threadIdx%x-1) + 1
    b(i) = a(j)
    return
    end subroutine

    attributes(global) subroutine withtex( a, b )
    real(8), device :: a(*), b(*)
    i = blockDim%x*(blockIdx%x-1) + threadIdx%x
    j = bitrev8(threadIdx%x-1) + 1
    b(i) = t(j)  ! This subroutine accesses a through the texture
    return
    end subroutine
end module memtests

program t
use cudafor
use memtests
real(8), device, target, allocatable :: da(:)
real(8), device, allocatable :: db(:)
integer, parameter :: n = 1024*1024
integer, parameter :: nthreads = 256
integer, parameter :: ntimes = 1000
type(cudaEvent) :: start, stop
real(8) b(n)

allocate(da(nthreads))
allocate(db(n))

istat = cudaEventCreate(start)
istat = cudaEventCreate(stop)

db = 100.0d0
da = (/ (dble(i),i=1,nthreads) /)

call without<<<n/nthreads, nthreads>>> (da, db)
istat = cudaEventRecord(start,0)
do j = 1, ntimes
  call without<<<n/nthreads, nthreads>>> (da, db)
end do
istat = cudaEventRecord(stop,0)
istat = cudaDeviceSynchronize()
istat = cudaEventElapsedTime(time1, start, stop)
time1 = time1 / (ntimes*1.0e3)
b = db
print *,sum(b)==(n*(nthreads+1)/2)

db = 100.0d0
t => da  ! assign the texture to da using f90 pointer assignment

call withtex<<<n/nthreads, nthreads>>> (da, db)
istat = cudaEventRecord(start,0)
do j = 1, ntimes
  call withtex<<<n/nthreads, nthreads>>> (da, db)
end do
istat = cudaEventRecord(stop,0)
istat = cudaDeviceSynchronize()
istat = cudaEventElapsedTime(time2, start, stop)
time2 = time2 / (ntimes*1.0e3)
b = db
print *,sum(b)==(n*(nthreads+1)/2)

print *,"Time with    textures",time2
print *,"Time without textures",time1
print *,"Speedup with textures",time1 / time2

deallocate(da)
deallocate(db)
end

The host attribute may be explicitly specified on the Subroutine or Function statement as follows:

attributes(host) subroutine sub(...) 
attributes(host) integer function func(...) 
integer attributes(host) function func(...) 

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 attribute may be explicitly specified on the Subroutine statement as follows:

attributes(global) subroutine sub(...)

Functions may not have the global attribute. A subroutine with the 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.

The device attribute may be explicitly specified on the Subroutine or Function statement as follows:

attributes(device) subroutine sub(...)
attributes(device) datatype function func(...)
datatype attributes(device) function func(...)

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.

A subroutine or function with the device or global attribute must satisfy the following restrictions:

• It may not be recursive, nor have the recursive prefix on the subprogram statement.

• It may not be pure or elemental, nor have the pure or elemental prefix on the subprogram statement.

• It may not contain another subprogram.

• It may not be contained in another subroutine or function.

For more information, refer to Device Code.

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.

real :: a(100)
attributes(device) :: a
real, device :: b(100)

These rules apply to device data:

• An allocatable device array dynamically allocates device global memory.

• 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.

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.

real :: a(100)
attributes(managed) :: a
real, managed :: b(100)
real, allocatable, managed :: c(:)
. . .
allocate(c(200))

These rules apply to managed data on the host:

• 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.

• 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 callingcudaSetStreamDefault.

• All subsequently allocated managed variables can also be associated with a given stream by callingcudaSetStreamDefault.

• 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.

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.

real :: p(:)
allocatable :: p
attributes(pinned) :: p
real, allocatable, pinned :: q(:)

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:

• As a named entity within a USE statement.

• As the source or destination in a data transfer assignment statement

• As an actual argument to another host subprogram

• As a dummy argument in a host subprogram

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.

attributes(global) subroutine sub(A, n,
integer, value :: n, nb
real, shared :: s(nb*blockdim%x,nb)

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:

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

    real, texture, pointer :: t(:)

2. 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.

3. 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.

    ! Vector add, s through device memory, t is through texture memory
    i = threadIdx%x + (blockIdx%x-1)*blockDim%x
    s(i) = s(i) + t(i)

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.

attributes(global) subroutine madd( a, b, n )
      real, dimension(n,n) :: a, b
      integer, value :: n

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.

real, allocatable, device :: b(:)
allocate(b(5024),stat=istat)
...
if(allocated(b)) deallocate(b)

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:

subroutine vfunc(a,c,n)
    real, device :: adev(n)
    real, device :: atmp(4)
    ...
end subroutine vfunc   ! adev and atmp are deallocated

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:

real, allocatable, device :: v(:)
istat = cudaMalloc(v, 100)
...
istat = cudaFree(v)

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

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.

real, allocatable, pinned :: p(:)
allocate(p(5000),stat=istat)
...
if(allocated(p)) deallocate(p)

To determine whether or not the allocation from page-locked memory was successful, an additional PINNED keyword is added to the allocate statement. It returns a logical success value.

logical plog
allocate(p(5000), stat=istat, pinned=plog)
if (.not. plog) then
. . .

Managed arrays may or may not have the allocatable attribute. These arrays are all dynamically allocated just as device arrays are.

real, allocatable, managed :: b(:)
allocate(b(5024),stat=istat)
...
if(allocated(b)) deallocate(b)

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:

subroutine vfunc(a,c,n)
    real, managed :: aman(n)
    real, managed :: atmp(4)
    ...
end subroutine vfunc   ! aman and atmp are deallocated

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

use cudafor
real, allocatable, managed :: v(:)
istat = cudaMallocManaged(v, 100, cudaMemAttachHost)
...
istat = cudaFree(v)

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

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:

real, device :: wrk(1024)
real cur(512)
istat = cudaMemcpy(wrk, cur, 512)

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.

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.

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

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

Several CUDA Fortran read-only predefined variables are available in device code. They are declared as follows:

type(dim3) :: threadidx, blockdim, blockidx, griddim
integer(4), parameter :: warpsize = 32

• 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.

This section lists the Fortran intrinsic functions allowed in device subprograms.

This section describes the new intrinsic functions and subroutines supported in device subprograms.

void syncthreads()

integer syncthreads_and(int_value)

integer syncthreads_count(int_value)

integer syncthreads_or(int_value)

void threadfence()

void threadfence_block()

void threadfence_system()

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

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

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

unsigned integer ballot(int_value)

if( ballot(int_value) ) allneg = .true.

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.

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

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

if (mem == comp) then
mem = val
endif

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.

PGI 14.1 enables CUDA Fortran device code to access compute capability 3.x shuffle functions. These functions enable access to variables between threads within a warp, referred to as lanes. In CUDA Fortran, lanes use Fortran's 1-based numbering scheme.

    integer(4) function __shfl(var, srcLane, width)
       integer(4) var, srcLane
       integer(4), optional :: width

    real(4) function __shfl(var, srcLane, width)
       real(4) :: var
       integer(4) :: srcLane
       integer(4), optional :: width

    real(8) function __shfl(var, srcLane, width)
       real(8) :: var
       integer(4) :: srcLane
       integer(4), optional :: width

    integer(4) function __shfl_up(var, delta, width)
       integer(4) var, delta
       integer(4), optional :: width

    real(4) function __shfl_up(var, delta, width)
       real(4) :: var
       integer(4) :: delta
       integer(4), optional :: width

    real(8) function __shfl_up(var, delta, width)
       real(8) :: var
       integer(4) :: delta
       integer(4), optional :: width

    integer(4) function __shfl_down(var, delta, width)
       integer(4) var, delta
       integer(4), optional :: width

    real(4) function __shfl_down(var, delta, width)
       real(4) :: var
       integer(4) :: delta
       integer(4), optional :: width

    real(8) function __shfl_down(var, delta, width)
       real(8) :: var
       integer(4) :: delta
       integer(4), optional :: width

    integer(4) function __shfl_xor(var, laneMask, width)
       integer(4) var, laneMask
       integer(4), optional :: width

    real(4) function __shfl_xor(var, laneMask, width)
       real(4) :: var
       integer(4) :: laneMask
       integer(4), optional :: width

    real(8) function __shfl_xor(var, laneMask, width)
       real(8) :: var
       integer(4) :: laneMask
       integer(4), optional :: width

        j = . . .
        k = __shfl_xor(j,1);  j = j + k
        k = __shfl_xor(j,2);  j = j + k
        k = __shfl_xor(j,4);  j = j + k
        k = __shfl_xor(j,8);  j = j + k
        k = __shfl_xor(j,16); j = j + k

This section lists restrictions on statements and features that can appear in device subprograms.

• Recursive subroutines and functions are not allowed.

• PAUSE statements are not allowed.

• Most Input/Output statements are not allowed at all: READ, FORMAT, NAMELIST, OPEN, CLOSE, BACKSPACE, REWIND, ENDFILE, INQUIRE.

• List-directed PRINT and WRITE statements to the default unit may be used when compiling for compute capability 2.0 and higher; all other uses of PRINT and WRITE are disallowed.

• Alternate return specifications are not allowed.

• ENTRY statements are not allowed.

• Floating point exception handling is not supported.

• Fortran intrinsic functions not listed in Section 3.6.3 are not supported.

• Cray pointers are not supported.

A call to sizeof(A), where A is a variable or expression, returns the number of bytes required to hold the value of A.

integer(kind=4) :: i, j
j = sizeof(i)    ! this assigns the value 4 to j

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:

   program multidimred
   use cudafor
   real(8), managed :: a(5,5,5,5,5)
   real(8), managed :: b(5,5,5,5)
   real(8) :: c
   call random_number(a)
   do idim = 1, 5
     b = sum(a, dim=idim)
     c = max(maxval(b), c)
   end do
   print *,"Max along any dimension",c
   end program

Array slices are also supported. This may run less efficiently on the GPU, but is very powerful nonetheless, and useful for debugging:

   real(4), managed :: a(n,m)
   reslt(ix) = sum(a(2:n-1,:))
   reslt(ix) = sum(a(:,3:m-2))
   reslt(ix) = sum(a(n2:n,m2:m))
   reslt(ix) = sum(a(1:n3,1:m3))
   reslt(ix) = sum(a(n2:n3,m2:m3))

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:

   program reductionRename
   use cudafor, gpusum => sum
   implicit none
   integer, managed :: m(3000)
   integer :: istat
   m = 1
   istat = cudaDeviceSynchronize()
   write(*,*) sum(m)    ! executes on host
   write(*,*) gpusum(m) ! executes on device
   end program

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

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

To access this module, 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.

CUDA Built-in Routines lists the CUDA built-in routines that are available:

PGI also provides a device module which provides interfaces to standard libm functions which are not in the Fortran intrinsic library.

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

use libm

These interfaces are defined in the libm device module:

On NVIDIA GPUs which support CUDA Compute Capability 3.5 and above, PGI provides a device module which provides interfaces to the device CUBLAS library.

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

use cublas_device

Here is a simple example of using the CUBLAS device library:

module mblasd
contains
  attributes(global) subroutine testblas( a, b, n )
  use cublas_device
  real, device :: a(*), b(*)
  type(cublasHandle) :: h
  integer, value :: n
  j = cublasCreate_v2(h)
  i = threadIdx%x
  if (i.eq.1) then
    j = cublasSswap_v2(h,n,a,1,b,1)
  end if
  j = cublasDestroy_v2(h)
  return
  end subroutine
end module mblasd

! Compile with "pgfortran -Mcuda=cc35 cublasdev.cuf -lcublas_device"
program cublasdev
use mblasd
integer, parameter :: n = 1000
real, managed :: a(n), b(n)
a = 1.0; b = 2.0
call testblas<<<1,1>>> (a,b,n)
istat = cudaDeviceSynchronize()
print *,all(a.eq.2*b)
end

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.

      call saxpy(n, a, x, incx, y, incy)

  where the arguments x and y have the device attribute.

• Portable legacy CUBLAS interfaces which interface directly with CUBLAS versions < 4.0, i.e.

      call cublasSaxpy(n, a, x, incx, y, incy)

  where the arguments x and y must have the device attribute.

• 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.

CUDA 3.2 helper functions defined in the cublas module:

integer function cublasInit()
integer function cublasShutdown()
integer function cublasGetError()
integer function cublasAlloc(n, elemsize, devptr)
integer function cublasFree(devptr)
integer function cublasSetVector(n, elemsize, x, incx, y, incy)
integer function cublasGetVector(n, elemsize, x, incx, y, incy)
integer function cublasSetMatrix(rows, cols, elemsize, a, lda, b, ldb)
integer function cublasGetMatrix(rows, cols, elemsize, a, lda, b, ldb)
integer function cublasSetKernelStream(stream)
integer function cublasSetVectorAsync(n, elemsize, x, incx, y, incy, stream)
integer function cublasGetVectorAsync(n, elemsize, x, incx, y, incy, stream)
integer function cublasSetMatrixAsync(rows, cols, elemsize, a, lda, b, ldb, stream)
integer function cublasGetMatrixAsync(rows, cols, elemsize, a, lda, b, ldb, stream)

Additional CUDA 4.0 helper functions defined in the cublas module:

integer function cublasCreate(handle)
integer function cublasDestroy(handle)
integer function cublasGetVersion(handle, version)
integer function cublasSetStream(handle, stream)
integer function cublasGetStream(handle, stream)
integer function cublasGetPointerMode(handle, mode)
integer function cublasSetPointerMode(handle, mode)

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.

use cufft

Here is an example of some code which uses the cufft interfaces:

program cufft2dTest
  use cufft
  integer, parameter :: n=450
  complex :: a(n,n),b(n,n)
  complex, device :: a_d(n,n), b_d(n,n)
  integer :: plan, ierr
  real, dimension(3) :: res, exp

  a  = 1; a_d  = a

  ierr = cufftPlan2D(plan,n,n,CUFFT_C2C)
  ierr = ierr + cufftExecC2C(plan,a_d,b_d,CUFFT_FORWARD)
  ierr = ierr + cufftExecC2C(plan,b_d,b_d,CUFFT_INVERSE)

  b = b_d
  res(1) = maxval(abs(a-b/(n*n)))
  print *,'Max error C2C: ', res(1)

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:

! Compile with "pgfortran testLevel3.cuf -lcusparse"
program testLevel3
  use cudafor
  use cusparse

  implicit none

  integer, parameter :: nd = 20 ! # rows/cols in dense matrix

  type(cusparseHandle) :: h
  type(cusparseMatDescr) :: descrA
  type(cusparseSolveAnalysisInfo) :: saInfo
  integer :: status, version, mode, i

  ! D-data
  ! dense
  real(8) :: DAde(nd,nd), DBde(nd,nd), DCde(nd,nd), DmaxErr
  real(8), device :: DAde_d(nd,nd), DBde_d(nd,nd), DCde_d(nd,nd)
  ! csr
  real(8) :: csrValDA(nd)
  real(8), device :: csrValDA_d(nd)
  real(8) :: Dalpha, Dbeta
  real(8), device :: Dalpha_d, Dbeta_d

  ! integer data common to all data types
  integer :: nnz
  integer :: nnzPerRowA(nd), csrRowPtrA(nd+1), csrColIndA(nd)
  integer, device :: nnzPerRowA_d(nd), csrRowPtrA_d(nd+1), csrColIndA_d(nd)

  ! initalization

  status = cusparseCreate(h)
  status = cusparseGetVersion(h, version)
  write(*,*) '... version:', version

  status = cusparseCreateMatDescr(descrA)
  status = cusparseSetMatType(descrA, CUSPARSE_MATRIX_TYPE_GENERAL)
  status = cusparseSetMatIndexBase(descrA, CUSPARSE_INDEX_BASE_ONE)
  status = cusparseCreateSolveAnalysisInfo(saInfo)

  ! Initialize matrix (Identity)

  DAde = 0.0
  i = 1, nd
      DAde(i,i) = 1.0
  end do
  DAde_d = DAde
  call random_number(DBde)
  DBde_d = DBde

  ! convert from dense to csr
  status = cusparseDnnz_v2(h, CUSPARSE_DIRECTION_ROW, nd, nd, descrA, &
      DAde_d, nd, nnzPerRowA_d, nnz)
  status = cusparseDdense2csr(h, nd, nd, descrA, DAde_d, nd, nnzPerRowA_d, &
      csrValDA_d, csrRowPtrA_d, csrColIndA_d)

  ! csrmm HPM
  Dalpha = 1.0
  Dbeta = 0.0
  status = cusparseDcsrmm(h, CUSPARSE_OPERATION_NON_TRANSPOSE, nd, nd, nd, &
      nnz, Dalpha, descrA, csrValDA_d, csrRowPtrA_d, csrColIndA_d, DBde_d, &
      nd, Dbeta, DCde_d, nd)
  if (status /= CUSPARSE_STATUS_SUCCESS) write (*,*) 'CSRMM Error:',status

  DCde = DCde_d
  DmaxErr = maxval(abs(DCde-DBde))

  status = cusparseDestroy(h)
  write(*,*) 'cusparseDestroy', status, DmaxErr

end program testLevel3

integer function cudaChooseDevice ( devnum, prop )
     integer, intent(out) :: devnum
     type(cudadeviceprop), intent(in) :: prop

cudaChooseDevice assigns the device number that best matches the properties given in prop to its first argument.

integer function cudaDeviceGetCacheConfig ( cacheconfig )
     integer, intent(out) :: cacheconfig

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.

integer function cudaDeviceGetLimit( val, limit )
   integer(kind=cuda_count_kind) :: val
   integer :: limit

cudaGetDeviceGetLimit returns in val the current size of limit. Current possible limit arguments are cudaLimitStackSize, cudaLimitPrintfSize, and cudaLimitMallocHeapSize.

cudaGetDeviceGetLimit is available in device code starting in CUDA 5.0.

integer function cudaDeviceGetSharedMemConfig ( config )
integer, intent(out) :: config

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.

integer function cudaDeviceGetStreamPriorityRange ( leastpriority, greatestpriority )
integer, intent(out) :: leastpriority, greatestpriority

cudaDeviceGetStreamPriorityRange returns the numerical values that correspond to the least and greatest stream priorities for the current context and device.

integer function cudaDeviceReset()

cudaDeviceReset resets the current device attached to the current process.

integer function cudaDeviceSetCacheConfig ( cacheconfig )
     integer, intent(in) :: cacheconfig

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

integer function cudaDeviceSetLimit( limit, val )
  integer :: limit

  integer(kind=cuda_count_kind) :: val

cudaGetDeviceSetLimit sets the limit of the current device to val. Current possible limit arguments are cudaLimitStackSize, cudaLimitPrintfSize, and cudaLimitMallocHeapSize.

integer function cudaDeviceSetSharedMemConfig ( config )
	integer, intent(in) :: config

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.

integer function cudaDeviceSynchronize()

cudaDeviceSynchronize blocks the current device until all preceding requested tasks have completed.

cudaDeviceSynchronize is available in device code starting in CUDA 5.0.

integer function cudaGetDevice( devnum )
	integer, intent(out) :: devnum

cudaGetDevice assigns the device number associated with this host thread to its first argument.

cudaGetDevice is available in device code starting in CUDA 5.0.

integer function cudaGetDeviceCount( numdev )
	integer, intent(out) :: numdev

cudaGetDeviceCount assigns the number of available devices to its first argument.

cudaGetDeviceCount is available in device code starting in CUDA 5.0.

integer function cudaGetDeviceProperties( prop, devnum )
     type(cudadeviceprop), intent(out) :: prop
     integer, intent(in) :: devnum

cudaGetDeviceProperties returns the properties of a given device.

cudaGetDeviceProperties is available in device code starting in CUDA 5.0.

integer function cudaSetDevice( devnum )
	integer, intent(in) :: devnum

cudaSetDevice selects the device to associate with this host thread.

integer function cudaSetDevice( flags )
	integer, intent(in) :: flags

cudaSetDeviceFlags records how the CUDA runtime interacts with this host thread.

integer function cudaSetValidDevices( devices, numdev )
	integer :: numdev, devices(numdev)

cudaSetValidDevices sets a list of valid devices for CUDA execution in priority order as specified in the devices array.

integer function cudaThreadExit()

cudaThreadExit explicitly cleans up all runtime-related CUDA resources associated with the host thread. Any subsequent CUDA calls or operations will reinitialize the runtime.

Calling cudaThreadExit is optional; it is implicitly called when the host thread exits.

integer function cudaThreadSynchronize()

cudaThreadSynchronize blocks execution of the host subprogram until all preceding kernels and operations are complete. It may return an error condition if one of the preceding operations fails.

function cudaGetErrorString( errcode )
    integer, intent(in) :: errcode
    character*(*) :: cudaGetErrorString

cudaGetErrorString returns the message string associated with the given error code.

integer function cudaGetLastError()

cudaGetLastError returns the error code that was most recently returned from any runtime call in this host thread.

integer function cudaPeekAtLastError()

cudaPeekAtLastError returns the last error code that has been produced by the CUDA runtime without resetting the error code to cudaSuccess like cudaGetLastError.

integer(kind=cuda_stream_kind) function cudaforGetDefaultStream( devptr )

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().

integer function cudaforSetDefaultStream( devptr, stream )
    integer(kind=cuda_stream_kind), intent(in) :: stream

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++.

integer function cudaStreamAttachMemAsync( stream, devptr, length, flags )
	integer(kind=cuda_stream_kind), intent(in) :: stream
	integer(kind=cuda_count_kind), optional, intent(in) :: length
	integer, optional, intent(in) :: flags
	

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.

cudaStreamAttachMemAsync is available starting in CUDA 6.0.

integer function cudaStreamCreate( stream )
	integer(kind=cuda_stream_kind), intent(out) :: stream

cudaStreamCreate creates an asynchronous stream and assigns its identifier to its first argument.

integer function cudaStreamCreateWithFlags( stream, flags )
	integer(kind=cuda_stream_kind), intent(out) :: stream
	integer, intent(in) :: flags

cudaStreamCreateWithFlags creates an asynchronous stream and assigns its identifier to its first argument. Valid values for flags are cudaStreamDefault or cudaStreamNonBlocking.

cudaStreamCreateWithFlags is available in device code starting in CUDA 5.0.

integer function cudaStreamCreateWithPriority( stream, flags, priority )
	integer(kind=cuda_stream_kind), intent(out) :: stream
	integer, intent(in) :: flags, priority

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.

integer function cudaStreamDestroy( stream )
	integer(kind=cuda_stream_kind), intent(in) :: stream

cudaStreamDestroy releases any resources associated with the given stream.

cudaStreamDestroy is available in device code starting in CUDA 5.0.

integer function cudaStreamGetPriority( stream, priority )
	integer(kind=cuda_stream_kind), intent(in) :: stream
	integer, intent(out) :: priority 

cudaStreamGetPriority queries and returns the priority of the given stream in priority.

integer function cudaStreamQuery( stream )
	integer(kind=cuda_stream_kind), intent(in) :: 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.

integer function cudaStreamSynchronize( stream )
	integer(kind=cuda_stream_kind), intent(in) :: stream

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.

integer function cudaStreamWaitEvent( stream, event, flags )
    integer(kind=cuda_stream_kind) :: stream
    type(cudaEvent), intent(in) :: event
    integer :: flags

cudaStreamWaitEvent blocks execution on all work submitted on the stream until the event reports completion.

cudaStreamWaitEvent is available in device code starting in CUDA 5.0.

integer function cudaEventCreate( event )
	type(cudaEvent), intent(out) :: event

cudaEventCreate creates an event object and assigns the event identifier to its first argument

integer function cudaEventCreateWithFlags( event, flags )
    type(cudaEvent), intent(out) :: event
    integer :: flags

cudaEventCreateWithFlags creates an event object with the specified flags. Current flags supported are cudaEventDefault, cudaEventBlockingSync, and cudaEventDisableTiming.

cudaEventCreateWithFlags is available in device code starting in CUDA 5.0.

integer function cudaEventDestroy( event )
	type(cudaEvent), intent(in) :: event

cudaEventDestroy destroys the resources associated with an event object.

cudaEventDestroy is available in device code starting in CUDA 5.0.

integer function cudaEventElapsedTime( time, start, end)
    real :: time
    type(cudaEvent), intent() :: start, end

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.

integer function cudaEventQuery( event )
    type(cudaEvent), intent(in) :: event

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.

integer function cudaEventRecord( event, stream )
    type(cudaEvent), intent(in) :: event
    integer, intent(in) :: stream

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.

cudaEventRecord is available in device code starting in CUDA 5.0.

integer function cudaEventSynchronize( event )
	type(cudaEvent), intent(in) :: event

cudaEventSynchronize blocks until the event has been recorded. It returns a value of cudaErrorInvalidValue if cudaEventRecord has not been called for this event.

integer function cudaFuncGetAttributes( attr, func )
    type(cudaFuncAttributes), intent(out) :: attr
    character*(*) :: func

cudaFuncGetAttributes gets the attributes for the function named by the func argument, which must be a global function.

cudaFuncGetAttributes is available in device code starting in CUDA 5.0.

integer function cudaFuncSetCacheConfig( func, cacheconfig )
    character*(*) :: func
    integer :: cacheconfig

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.

integer function cudaFuncSetSharedMemConfig( func, cacheconfig )
character*(*) :: func
integer :: cacheconfig

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

integer function cudaSetDoubleForDevice( d )
    real(8) :: d

cudaSetDoubleForDevice sets the argument d to an internal representation suitable for devices which do not support double precision arithmetic.

integer function cudaSetDoubleForHost( d )
	real(8) :: d

cudaSetDoubleForHost sets the argument d from an internal representation on devices which do not support double precision arithmetic to the normal host representation.

integer function cudaFree(devptr)

cudaFree deallocates data on the device. devptr may be any allocatable device array of a supported type specified in Device Code Intrinsic Datatypes. Or, devptr may be of TYPE(C_DEVPTR).

cudaFree is available in device code starting in CUDA 5.0.

integer function cudaFreeArray(carray)
	type(cudaArrayPtr) :: carray

cudaFreeArray frees an array that was allocated on the device.

integer function cudaFreeHost(hostptr)
    type(C_PTR) :: hostptr

cudaFreeHost deallocates pinned memory on the host allocated with cudaMalloHost.

integer function cudaGetSymbolAddress(devptr, symbol)
    type(C_DEVPTR) :: devptr
    type(c_ptr) :: symbol

cudaGetSymbolAddress returns in the devptr argument the address of symbol on the device. A symbol can be set to an external device name via a character string.

The following code sequence initializes a global device array ‘vx’ from a CUDA C kernel:

type(c_ptr) :: csvx
type(c_devptr) :: cdvx
real, allocatable, device :: vx(:)
csvx = ‘vx’
Istat = cudaGetSymbolAddress(cdvx, csvx)
Call c_f_pointer(cdvx, vx, 100)
Vx = 0.0

integer function cudaGetSymbolSize(size, symbol)
    integer :: size
    type(c_ptr) :: symbol

cudaGetSymbolSize sets the variable size to the size of a device area in global or constant memory space referenced by the symbol.

integer function cudaHostAlloc(hostptr, size, flags)
    type(C_PTR) :: hostptr
    integer :: size, flags

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.

integer function cudaHostGetDevicePointer(devptr, hostptr, flags)
    type(C_DEVPTR) :: devptr
    type(C_PTR) :: hostptr
    integer :: flags

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.

integer function cudaHostGetFlags(flags, hostptr)
    integer :: flags
    type(C_PTR) :: hostptr

cudaHostGetFlags returns the flags associated with a host pointer.

integer function cudaHostRegister(hostptr, count, flags)
    integer :: flags
    type(C_PTR) :: hostptr

cudaHostRegister page-locks the memory associated with the host pointer and of size provided by the count argument, according to the flags argument.

integer function cudaHostRegister(hostptr)
	type(C_PTR) :: hostptr

cudaHostUnregister unmaps the memory associated with the host pointer and makes it page-able again. The argument hostptr must be the same as was used with cudaHostRegister.

integer function cudaMalloc(devptr, count)

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.

cudaMalloc is available in device code starting in CUDA 5.0.

integer function cudaMallocArray(carray, cdesc, width, height)
    type(cudaArrayPtr) :: carray
    type(cudaChannelFormatDesc) :: cdesc
    integer :: width, height

cudaMallocArray allocates a data array on the device.

integer function cudaMallocManaged(devptr, count, flags)

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.

cudaMallocManaged is available starting in CUDA 6.0.

integer function cudaMallocPitch(devptr, pitch, width, height)

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.

integer function cudaMalloc3D(pitchptr, cext)
    type(cudaPitchedPtr), intent(out) :: pitchptr
    type(cudaExtent), intent(in) :: cext

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.

integer function cudaMalloc3DArray(carray, cdesc, cext)
    type(cudaArrayPtr) :: carray
    type(cudaChannelFormatDesc) :: cdesc
    type(cudaExtent) :: cext

cudaMalloc3DArray allocates array data on the device.

integer function cudaMemAdvise(devptr, count, advice, device)

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.

integer function cudaMemcpy(dst, src, count, kdir)

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.

cudaMemcpy is available in device code starting in CUDA 5.0.

integer function cudaMemcpyArrayToArray(dsta, dstx, dsty,
                        srca, srcx, srcy, count, kdir)
    type(cudaArrayPtr) :: dsta, srca
    integer :: dstx, dsty, srcx, srcy, count, kdir

cudaMemcpyArrayToArray copies array data to and from the device.

integer function cudaMemcpyAsync(dst, src, count, kdir, stream)

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.

cudaMemcpyAsync is available in device code starting in CUDA 5.0.

integer function cudaMemcpyFromArray(dst, srca, srcx, srcy, count, kdir)
    type(cudaArrayPtr) :: srca
    integer :: dstx, dsty, count, kdir

cudaMemcpyFromArray copies array data to and from the device.

integer function cudaMemcpyFromSymbol(dst, symbol, count, offset, kdir, stream)
    type(c_ptr) :: symbol
    integer :: count, offset, kdir
    integer, optional :: stream

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.

integer function cudaMemcpyFromSymbolAsync(dst, symbol, count, offset, kdir, stream)
    type(c_ptr) :: symbol
    integer :: count, offset, kdir
    integer, optional :: stream

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.

integer function cudaMemcpyPeer(dst, dstdev, src, srcdev, count)

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.

integer function cudaMemcpyPeerAsync(dst, dstdev, src, srcdev, count, stream)

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.

integer function cudaMemcpyToArray(dsta, dstx, dsty, src, count, kdir)
    type(cudaArrayPtr) :: dsta
    integer :: dstx, dsty, count, kdir

cudaMemcpyToArray copies array data to and from the device.

integer function cudaMemcpyToSymbol(symbol, src, count, offset, kdir)
   type(c_ptr) :: symbol
   integer :: count, offset, kdir

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.

integer function cudaMemcpyToSymbolAsync(symbol, src, count, offset, kdir, stream)
    type(c_ptr) :: symbol
    integer :: count, offset, kdir
    integer, optional :: stream

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.

integer function cudaMemcpy2D(dst, dpitch, src, spitch, width, height, kdir)

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 is available in device code starting in CUDA 5.0.

integer function cudaMemcpy2DArrayToArray(dsta, dstx, dsty,
                     srca, srcx, srcy, width, height, kdir)
    type(cudaArrayPtr) :: dsta, srca
    integer :: dstx, dsty, srcx, srcy, width, height, kdir

cudaMemcpy2DArrayToArray copies array data to and from the device.

integer function cudaMemcpy2DAsync(dst, dpitch, src, spitch, width,
              height, kdir, stream)

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.

cudaMemcpy2DAsync is available in device code starting in CUDA 5.0.

integer function cudaMemcpy2DFromArray(dst, dpitch, srca, srcx, srcy, 
                               width, height, kdir) 
    type(cudaArrayPtr) :: srca 
    integer :: dpitch, srcx, srcy, width, height, kdir

cudaMemcpy2DFromArray copies array data to and from the device.

integer function cudaMemcpy2DToArray(dsta, dstx, dsty, src, 
                     spitch, width, height, kdir)
    type(cudaArrayPtr) :: dsta
    integer :: dstx, dsty, spitch, width, height, kdir

cudaMemcpy2DToArray copies array data to and from the device.

integer function cudaMemcpy3D(p)
	type(cudaMemcpy3DParms) :: p

cudaMemcpy3D copies elements from one 3D array to another as specified by the data held in the derived type p.

integer function cudaMemcpy3D(p, stream)
    type(cudaMemcpy3DParms) :: p
    integer :: stream

cudaMemcpy3DAsync copies elements from one 3D array to another as specified by the data held in the derived type p.

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

integer function cudaMemGetInfo( free, total )
    integer(kind=cuda_count_kind) :: free, total

cudaMemGetInfo returns the amount of free and total memory available for allocation on the device. The returned values units are in bytes.

integer function cudaMemPrefetchAsync(devptr, count, device, stream)

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.

integer function cudaMemset(devptr, value, count)

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.

integer function cudaMemset2D(devptr, pitch, value, width, height)

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.

integer function cudaMemset3D(pitchptr, value, cext)
    type(cudaPitchedPtr) :: pitchptr
    integer :: value 
    type(cudaExtent) :: cext

cudaMemset3D sets elements of an array, the extents in each dimension specified by cext, which was allocated with cudaMalloc3D to a specified value.

integer function cudaDeviceCanAccessPeer( canAccessPeer, device, peerDevice )
	integer :: canAccessPeer, device, peerDevice

cudaDeviceCanAccessPeer returns in canAccessPeer the value 1 if the device argument can access memory in the device specified by the peerDevice argument.

integer function cudaDeviceDisablePeerAccess ( peerDevice )
	integer :: peerDevice

cudaDeviceDisablePeerAccess disables the ability to access memory on the device specified by the peerDevice argument by the current device.

integer function cudaDeviceEnablePeerAccess ( peerDevice, flags )
	integer :: peerDevice, flags

cudaDeviceEnablePeerAccess enables the ability to access memory on the device specified by the peerDevice argument by the current device. Currently, flags must be zero.

integer function cudaPointerGetAttributes( attr, ptr )
    type(cudaPointerAttributes), intent(out) :: ptr

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.

integer function cudaDriverGetVersion(iversion)
	integer :: iversion

cudaDriverGetVersion returns the version number of the installed CUDA driver as iversion. If no driver is installed, then it returns 0 as iversion.

This function automatically returns cudaErrorInvalidValue if the iversion argument is NULL.

integer function cudaRuntimeGetVersion(iversion)
	integer :: iversion

cudaRuntimeGetVersion returns the version number of the installed CUDA Runtime as iversion.

This function automatically returns cudaErrorInvalidValue if the iversion argument is NULL.

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