Visible to Intel only — GUID: GUID-DB415E02-7B1A-4BDE-8918-724BD40E5326
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
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
Visible to Intel only — GUID: GUID-DB415E02-7B1A-4BDE-8918-724BD40E5326
Making Better Use of OpenMP Constructs
Reduce Synchronizations Using nowait
If appropriate, use the nowait clause on the target construct to reduce synchronizations.
By default, there is an implicit barrier at the end of a target region, which ensures that the host thread that encountered the target construct cannot continue until the target region is complete.
Adding the nowait clause on the target construct eliminates this implicit barrier, so the host thread that encountered the target construct can continue even if the target region is not complete. This allows the target region to execute asynchronously on the device without requiring the host thread to idly wait for the target region to complete.
Consider the following example, which computes the product of two vectors, v1 and v2, in a parallel region (line 48). Half of the computations are performed on the host by the team of threads executing the parallel region. The other half of the computations are performed on the device. The master thread of the team launches a target region to do the computations on the device.
By default, the master thread of the team has to wait for the target region to complete before proceeding and participating in the computations (worksharing for loop) on the host.
/*
* This test is taken from OpenMP API 5.0.1 Examples (June 2020)
* https://www.openmp.org/wp-content/uploads/openmp-examples-5-0-1.pdf
* (4.13.2 nowait Clause on target Construct)
*/
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <omp.h>
#define N 100000 // N must be even
void init(int n, float *v1, float *v2) {
int i;
for(i=0; i<n; i++){
v1[i] = i * 0.25;
v2[i] = i - 1.25;
}
}
int main() {
int i, n=N;
float v1[N],v2[N],vxv[N];
double start,end; // timers
init(n, v1, v2);
/* Dummy parallel and target regions, so as not to measure startup
time. */
#pragma omp parallel
{
#pragma omp master
#pragma omp target
{;}
}
start=omp_get_wtime();
#pragma omp parallel
{
#pragma omp master
#pragma omp target teams distribute parallel for \
map(to: v1[0:n/2]) \
map(to: v2[0:n/2]) \
map(from: vxv[0:n/2])
for(i=0; i<n/2; i++){
vxv[i] = v1[i]*v2[i];
}
/* Master thread will wait for target region to be completed
before proceeding beyond this point. */
#pragma omp for
for(i=n/2; i<n; i++) {
vxv[i] = v1[i]*v2[i];
}
/* Implicit barrier at end of worksharing for. */
}
end=omp_get_wtime();
printf("vxv[0]=%f, vxv[n-1]=%f, time=%lf\n", vxv[0], vxv[n-1], end-start);
return 0;
}
Compilation command:
icpx -fiopenmp -fopenmp-targets=spir64 test_target_no_nowait.cpp
Run command:
OMP_TARGET_OFFLOAD=MANDATORY ZE_AFFINITY_MASK=0.0 LIBOMPTARGET_DEBUG=1 ./a.out
Performance could be improved if a nowait clause is specified on the target construct, so the master thread does not have to wait for the target region to complete and can proceed to work on the worksharing for loop. The target region is guaranteed to complete by the synchronization in the implicit barrier at the end of the worksharing for loop.
/*
* This test is taken from OpenMP API 5.0.1 Examples (June 2020)
* https://www.openmp.org/wp-content/uploads/openmp-examples-5-0-1.pdf
* (4.13.2 nowait Clause on target Construct)
*/
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <omp.h>
#define N 100000 // N must be even
void init(int n, float *v1, float *v2) {
int i;
for(i=0; i<n; i++){
v1[i] = i * 0.25;
v2[i] = i - 1.25;
}
}
int main() {
int i, n=N;
float v1[N],v2[N],vxv[N];
double start,end; // timers
init(n, v1,v2);
/* Dummy parallel and target (nowait) regions, so as not to measure
startup time. */
#pragma omp parallel
{
#pragma omp master
#pragma omp target nowait
{;}
}
start=omp_get_wtime();
#pragma omp parallel
{
#pragma omp master
#pragma omp target teams distribute parallel for nowait \
map(to: v1[0:n/2]) \
map(to: v2[0:n/2]) \
map(from: vxv[0:n/2])
for(i=0; i<n/2; i++){
vxv[i] = v1[i]*v2[i];
}
#pragma omp for
for(i=n/2; i<n; i++) {
vxv[i] = v1[i]*v2[i];
}
/* Implicit barrier at end of worksharing for. Target region is
guaranteed to be completed by this point. */
}
end=omp_get_wtime();
printf("vxv[1]=%f, vxv[n-1]=%f, time=%lf\n", vxv[1], vxv[n-1], end-start);
return 0;
}
The performance of the two versions when running on one of our lab machines was as follows:
no nowait version : 0.008220 seconds nowait on target version : 0.002110 seconds
Fortran
The same nowait example shown above may be written in Fortran as follows.
!
! This test is from OpenMP API 5.0.1 Examples (June 2020)
! https://www.openmp.org/wp-content/uploads/openmp-examples-5-0-1.pdf
!(4.13.2 nowait Clause on target Construct)
!
subroutine init(n, v1, v2)
integer :: i, n
real :: v1(n), v2(n)
do i = 1, n
v1(i) = i * 0.25
v2(i) = i - 1.25
end do
end subroutine init
program test_target_nowait
use omp_lib
use iso_fortran_env
implicit none
integer, parameter :: NUM=100000 ! NUM must be even
real :: v1(NUM), v2(NUM), vxv(NUM)
integer :: n, i
real(kind=REAL64) :: start, end
n = NUM
call init(n, v1, v2)
! Dummy parallel and target (nowait) regions, so as not to measure
! startup time.
!$omp parallel
!$omp master
!$omp target nowait
!$omp end target
!$omp end master
!$omp end parallel
start=omp_get_wtime()
!$omp parallel
!$omp master
!$omp target teams distribute parallel do nowait &
!$omp& map(to: v1(1:n/2)) &
!$omp& map(to: v2(1:n/2)) &
!$omp& map(from: vxv(1:n/2))
do i = 1, n/2
vxv(i) = v1(i)*v2(i)
end do
!$omp end master
!$omp do
do i = n/2+1, n
vxv(i) = v1(i)*v2(i)
end do
!$omp end parallel
end=omp_get_wtime()
write(*,110) "vxv(1)=", vxv(1), ", vxv(n-1)=", vxv(n-1), ", time=", end-start
110 format (A, F10.6, A, F17.6, A, F10.6)
end program test_target_nowait