GPU621/To Be Announced

GPU621/DPS921 | Participants | Groups and Projects | Resources | Glossary

OpenMP Device Offloading

OpenMP 4.0/4.5 introduced support for heterogeneous systems such as accelerators and GPUs. The purpose of this overview is to demonstrate OpenMP's device constructs used for offloading data and code from a host device (Multicore CPU) to a target's device environment (GPU/Accelerator). We will demonstrate how to manage the device's data environment, parallelism and work-sharing. Review how data is mapped from the host data environment to the device data environment, and attempt to use different compilers that support OpenMP offloading such as LLVM/Clang or GCC.

Group Members

1. Elena Sakhnovitch

2. Nathan Olah

3. Yunseon Lee

Latest GPU specs (Yunseon)

AMD:

NVIDIA:

Means of parallelisation on GPUs

short introduction and advantages and disadvantages of:

CUDA (Yunseon)

OpenMP (Elena)

HIP (Elena)

OpenCL (Nathan) https://stackoverflow.com/questions/7263193/opencl-vs-openmp-performance#7263823

Instructions

How to set up compiler and target offloading for windows, on NVIDIA GPU: (Nathan)

How to set up compiler and target offloading for Linux on AMD GPU: (Elena)

Programming GPUs with OpenMP

Target Region

• The target region is the offloading construct in OpenMP.

int main() {
// This code executes on the host (CPU)

#pragma omp target
    // This code executes on the device

}

• An OpenMP program will begin executing on the host (CPU).
• When a target region is encountered the code that is within the target region will begin to execute on a device (GPU).

If no other construct is specified, for instance a construct to enable a parallelized region (#pragma omp parallel). By default, the code within the target region will execute sequentially. The target region does not express parallelism, it only expresses where the contained code is going to be executed on.

There is an implied synchronization between the host and the device at the end of a target region. At the end of a target region the host thread waits for the target region to finish execution and continues executing the next statements.

Mapping host and device data

• In order to access data inside the target region it must be mapped to the device.
• The host environment and device environment have separate memory.
• Data that has been mapped to the device from the host cannot access that data until the target region (Device) has completed its execution.

The map clause provides the ability to control a variable over a target region.

#pragma omp target map(map-type : list)

• list specifies the data variables to be mapped from the host data environment to the target's device environment.

• map-type is one of the types to, from, tofrom, or alloc.

to - copies the data to the device on execution.

from - copies the data to the host on exit.

tofrom - copies the data to the device on execution and back on exit.

alloc - allocated an uninitialized copy on the device (without copying from the host environment).

#pragma omp target map(to:A,B), map(tofrom:sum)
{
    for (int i = 0; i < N; i++) 
        sum += A[i] + B[i];
}

Declare Target

Calling functions within the scope of a target region.

#pragma omp declare target
    int combine(int a, int b);
#pragma omp end declare target

Sources

https://www.ibm.com/support/knowledgecenter/en/SSXVZZ_16.1.0/com.ibm.xlcpp161.lelinux.doc/compiler_ref/prag_omp_target.html https://www.ibm.com/support/knowledgecenter/en/SSXVZZ_16.1.0/com.ibm.xlcpp161.lelinux.doc/compiler_ref/prag_omp_declare_target.html