OpenMP: Fix memory access faults for non-USM offloads with 'nohost'
Corrects OpenMP's 'nohost' target handling to prevent memory access faults in non-USM offloads.
This change fixes a memory access fault occurring during OpenMP offloading when device_type(nohost) is used with non-Unified Shared Memory (non-USM) offloads. The issue manifested as a ‘Page not present’ error on AMD GPUs or an ‘illegal memory access’ on NVIDIA GPUs. The solution involves mapping the variable ‘x’ for the OpenMP ‘target’ construct, ensuring that data is correctly accessible on the device for non-USM scenarios.
In Details
The commit addresses a regression in OpenMP's device_type(nohost) handling for non-USM offloads, which caused memory access violations when targeting GPUs. By adding a map(x) clause to the OpenMP target construct in the test cases libgomp.c/target-device-type-1.c and libgomp.c/target-device-type-2.c, the test ensures that the variable 'x' is properly transferred to the device, resolving the runtime memory access errors experienced by both AMD and NVIDIA offloading implementations.
- non-USM
- Non-Unified Shared Memory. Refers to memory management in GPU offloading where the host and device have separate memory spaces. Data must be explicitly transferred between them, unlike Unified Shared Memory (USM) where a single memory address space is shared.
- map
- An OpenMP clause used with directives like
targetto specify how data should be transferred between the host and the device. For example,map(to:x)indicates that variable 'x' should be copied to the device. - offloading
- A technique in parallel computing where parts of a program are executed on a separate device, such as a GPU or a coprocessor, while the main part runs on the host CPU. OpenMP provides directives like
#pragma omp targetto facilitate offloading.