Visible to Intel only — GUID: GUID-9982BBC0-772C-4BEF-BEEE-DAB13B295EDA
Execution Model Overview
Thread Mapping and GPU Occupancy
Kernels
Using Libraries for GPU Offload
Host/Device Memory, Buffer and USM
Host/Device Coordination
Using Multiple Heterogeneous Devices
Compilation
OpenMP Offloading Tuning Guide
Multi-GPU, Multi-Stack and Multi-C-Slice Architecture and Programming
Level Zero
Performance Profiling and Analysis
Configuring GPU Device
Sub-Groups and SIMD Vectorization
Removing Conditional Checks
Registerization and Avoiding Register Spills
Porting Code with High Register Pressure to Intel® Max GPUs
Small Register Mode vs. Large Register Mode
Shared Local Memory
Pointer Aliasing and the Restrict Directive
Synchronization among Threads in a Kernel
Considerations for Selecting Work-Group Size
Prefetch
Reduction
Kernel Launch
Executing Multiple Kernels on the Device at the Same Time
Submitting Kernels to Multiple Queues
Avoiding Redundant Queue Constructions
Programming Intel® XMX Using SYCL Joint Matrix Extension
Doing I/O in the Kernel
Explicit Scaling on Multi-GPU, Multi-Stack, Multi-C-Slice in SYCL
Explicit Scaling Using Intel® oneAPI Math Kernel Library (oneMKL) in SYCL
Explicit Scaling on Multi-GPU, Multi-Stack and Multi-C-Slice in OpenMP
Explicit Scaling Using Intel® oneAPI Math Kernel Library (oneMKL) in OpenMP
Explicit Scaling Summary
Visible to Intel only — GUID: GUID-9982BBC0-772C-4BEF-BEEE-DAB13B295EDA
Clauses: is_device_ptr, use_device_ptr, has_device_addr, use_device_addr
The OpenMP clauses is_device_ptr, use_device_ptr, has_device_addr, and use_device_addr can be used to convey information about variables referenced in target, target data, or dispatch constructs. These clauses are described as follows.
is_device_ptr
The is_device_ptr clause appears on a target or dispatch directive. It indicates that the list items are device pointers. So each list item is privatized inside the construct and the new list item is initialized to the device address to which the original list item refers.
In C, each list item should be of type pointer or array.
In C++, each list item should be of type pointer, array, reference to pointer, or reference to array.
In Fortran, each list item should be of type C_PTR.
The following C/C++ example illustrates the use of the is_device_ptr clause. The omp_target_alloc_device routine allocates memory on the device and returns a device pointer for that memory which is saved in the host variable arr_device. On the target directive, we use the is_device_ptr(arr_device) clause to indicate that arr_device points to device memory. So inside the target construct arr_device is privatized and initialized to the device address to which arr_device refers.
#include <stdio.h>
#include <stdlib.h>
#include <stdint.h>
#include <math.h>
#include <omp.h>
#define N 100
int main(void)
{
int *arr_host = NULL;
int *arr_device = NULL;
arr_host = (int *) malloc(N * sizeof(int));
arr_device = (int *) omp_target_alloc_device(N * sizeof(int),
omp_get_default_device());
#pragma omp target is_device_ptr(arr_device) map(from: arr_host[0:N])
{
for (int i = 0; i < N; ++i) {
arr_device[i] = i;
arr_host[i] = arr_device[i];
}
}
printf ("%d, %d, %d \n", arr_host[0], arr_host[N/2], arr_host[N-1]);
}
use_device_ptr
The use_device_ptr clause appears on a target data directive. It indicates that each list item is a pointer to an object that has corresponding storage on the device or is accessible on the device.
If a list item is a pointer to an object that is mapped to the device, then references to the list item in the construct are converted to references to a device pointer that is local to the construct and that refers to the device address of the corresponding object.
If the list item does not point to a mapped object, it must contain a valid device address, and the list item references are converted to references to a local device pointer that refers to this device address.
Each list item must be a pointer for which the value is the address of an object that has corresponding storage in the device data environment or is accessible on the target device.
In C, each list item should be of type pointer or array.
In C++, each list item should be of type pointer, array, reference to pointer, or reference to array.
In Fortran, each list item should be of type C_PTR.
The following C/C++ example illustrates the use of the use_device_ptr clause. The omp_target_alloc_device routine is called three times to allocate memory on the device. The addresses of the memory allocated is saved in the pointer variables A, B, and C on the host. We use the use_device_ptr(A, B,C) clause on the target data directive to indicate that A, B, and C contain valid device addresses.
#include <stdio.h>
#include <stdlib.h>
#include <stdint.h>
#include <math.h>
#include <omp.h>
#define length 65536
int main(void)
{
int device_id = omp_get_default_device();
size_t bytes = length*sizeof(double);
double * __restrict A;
double * __restrict B;
double * __restrict C;
double scalar = 3.0;
double ar;
double br;
double cr;
double asum;
// Allocate arrays in device memory
A = (double *) omp_target_alloc_device(bytes, device_id);
if (A == NULL){
printf(" ERROR: Cannot allocate space for A using omp_target_alloc_device().\n");
exit(1);
}
B = (double *) omp_target_alloc_device(bytes, device_id);
if (B == NULL){
printf(" ERROR: Cannot allocate space for B using omp_target_alloc_device().\n");
exit(1);
}
C = (double *) omp_target_alloc_device(bytes, device_id);
if (C == NULL){
printf(" ERROR: Cannot allocate space for C using omp_target_alloc_device().\n");
exit(1);
}
#pragma omp target data use_device_ptr(A,B,C)
{
// Initialize the arrays
#pragma omp target teams distribute parallel for
for (size_t i=0; i<length; i++) {
A[i] = 2.0;
B[i] = 2.0;
C[i] = 0.0;
}
// Perform the computation
#pragma omp target teams distribute parallel for
for (size_t i=0; i<length; i++) {
C[i] += A[i] + scalar * B[i];
}
// Validate and output results
ar = 2.0;
br = 2.0;
cr = 0.0;
for (int i=0; i<length; i++) {
cr += ar + scalar * br;
}
asum = 0.0;
#pragma omp target teams distribute parallel for reduction(+:asum)
for (size_t i=0; i<length; i++) {
asum += fabs(C[i]);
}
} // end target data
omp_target_free(A, device_id);
omp_target_free(B, device_id);
omp_target_free(C, device_id);
double epsilon=1.e-8;
if (fabs(cr - asum)/asum > epsilon) {
printf("Failed Validation on output array\n"
" Expected checksum: %lf\n"
" Observed checksum: %lf\n"
"ERROR: solution did not validate\n", cr, asum);
return 1;
} else {
printf("Solution validates. Checksum = %lf\n", asum);
}
return 0;
}
has_device_addr
The has_device_addr clause appears on a target directive. It indicates that the list items already have valid device addresses, and therefore may be directly accessed from the device.
Each list item must have a valid device address for the device data environment. It can be on any type, including an array section.
The has_device_addr clause is especially useful in Fortran, because it can be used with list items of any type (not just C_PTR) to indicate that the list items have device addresses.
The following Fortran example illustrates the use of the has_device_addr clause. In the example, the three arrays A, B, and C are allocated on the device. When the arrays are referenced in a target region, we use the has_device_addr(A, B, C) clause to indicate that A, B, and C already have device addresses.
program main
use iso_fortran_env
use omp_lib
implicit none
integer, parameter :: iterations=1000
integer, parameter :: length=64*1024*1024
real(kind=REAL64), parameter :: epsilon=1.D-8
real(kind=REAL64), allocatable :: A(:)
real(kind=REAL64), allocatable :: B(:)
real(kind=REAL64), allocatable :: C(:)
real(kind=REAL64) :: scalar=3.0
real(kind=REAL64) :: ar, br, cr, asum
real(kind=REAL64) :: nstream_time, avgtime
integer :: i, iter
!
! Allocate arrays in device memory
!$omp allocators allocate(allocator(omp_target_device_mem_alloc): A)
allocate(A(length))
!$omp allocators allocate(allocator(omp_target_device_mem_alloc): B)
allocate(B(length))
!$omp allocators allocate(allocator(omp_target_device_mem_alloc): C)
allocate(C(length))
!
! Initialize the arrays
!$omp target teams distribute parallel do has_device_addr(A, B, C)
do i = 1, length
A(i) = 2.0
B(i) = 2.0
C(i) = 0.0
end do
!
! Perform the computation
nstream_time = omp_get_wtime()
do iter = 1, iterations
!$omp target teams distribute parallel do has_device_addr(A, B, C)
do i = 1, length
C(i) = C(i) + A(i) + scalar * B(i)
end do
end do
nstream_time = omp_get_wtime() - nstream_time
!
! Validate and output results
ar = 2.0
br = 2.0
cr = 0.0
do iter = 1, iterations
do i = 1, length
cr = cr + ar + scalar * br
end do
end do
asum = 0.0
!$omp target teams distribute parallel do reduction(+:asum) has_device_addr(C)
do i = 1, length
asum = asum + abs(C(i))
end do
if (abs(cr - asum)/asum > epsilon) then
print *, "Failed Validation on output array:", "Expected =", cr, "Observed =", asum
else
avgtime = nstream_time/iterations
print *, "Solution validates:", "Checksum =", asum, "Avg time (s) =", avgtime
endif
deallocate(A)
deallocate(B)
deallocate(C)
end program main
use_device_addr
The use_device_addr clause appears on a target data directive. It indicates that each list item already has corresponding storage on the device or is accessible on the device.
If a list item is mapped, then references to the list item in the construct are converted to references to the corresponding list item. If a list item is not mapped, it is assumed to be accessible on the device.
A list item may be an array section.
Just like has_device_addr, the use_device_addr clause is especially useful in Fortran, because it can be used with list items of any type (not just C_PTR) to indicate that the list items have device addresses.
The following Fortran example illustrates the use of the use_device_addr clause. In the example, array_d is mapped to the device with the alloc map-type, so storage is allocated for array_d on the device and no data transfer between the host and the device occurs. We use the use_device_addr(array_d) clause on the target data directive to indicate that array_d has corresponding storage on the device.
program target_use_device_addr
use omp_lib
use iso_fortran_env, only : real64
implicit none
integer, parameter :: N1 = 1024
real(kind=real64), parameter :: aval = real(42, real64)
real(kind=real64), allocatable :: array_d(:), array_h(:)
integer :: i,err
! Allocate host data
allocate(array_h(N1))
!$omp target data map (from:array_h(1:N1)) map(alloc:array_d(1:N1))
!$omp target data use_device_addr(array_d)
!$omp target
do i=1, N1
array_d(i) = aval
array_h(i) = array_d(i)
end do
!$omp end target
!$omp end target data
!$omp end target data
! Check result
write (*,*) array_h(1), array_h(N1)
if (any(array_h /= aval)) then
err = 1
else
err = 0
end if
deallocate(array_h)
if (err == 1) then
stop 1
else
stop 0
end if
end program target_use_device_addr
The following table summarizes the properties of the clauses described in this section.
Clause |
On which directive |
Type of list item |
Description |
|---|---|---|---|
is_device_ptr |
target, dispatch |
C/C++: Pointer, |
Indicates that list item is a device pointer (has valid device address). |
use_device_ptr |
target data |
C/C++: Pointer, |
Indicates that list item is a pointer to an object that has corresponding storage on device or is accessible on device. |
has_device_addr |
target |
Any type (may be array section) |
Indicates that list item has a valid device address. |
use_device_addr |
target data |
Any type (may be array section) |
Indicates that list item has corresponding storage on device or is accessible on the device. |