<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/127334>127334</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
#pragma omp requires unified_shared_memory does not seem to work
</td>
</tr>
<tr>
<th>Labels</th>
<td>
new issue
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
bschulz81
</td>
</tr>
</table>
<pre>
Unfortunately, openmp is very restricted in its
#teams distribute statement. it says:
"If a teams region is nested inside a target region, the corresponding target construct must not contain any statements, declarations or directives outside of the corresponding teams construct."
This leads to restrictions which are often not very practical.
One alternative, however, would be the unified shared memory of cuda/amd, which is available in openmp via
#pragma omp requires unified_shared_memory
yet this simple program fails. I start the compiler with this: -std=c++20 -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -lgomp
```
#include "omp.h"
#include <math.h>
#include <vector>
#include <numeric>
using namespace std;
#pragma omp requires unified_shared_memory
int main()
{
std::vector<double> v2(10) ;
std::iota (std::begin(v2), std::end(v2), 0);
std::vector<double> v(10,0) ;
double*d1=v.data();
double*d2=v2.data();
//no mapping needed in unified memory
//#pragma omp target enter data map(to: d1[0:10],d2[0:10])
#pragma omp target teams distribute
for(size_t i=1;i<10;i++)
{
d2[i]=d1[i]+1;
}
}
```
> "PluginInterface" error: Failure to synchronize stream (nil): Error in cuStreamSynchronize: an illegal memory access was encountered
> omptarget error: Consult https://openmp.llvm.org/design/Runtimes.html for debugging options.
> omptarget error: Source location information not present. Compile with -g or -gline-tables-only.
> omptarget fatal error 1: failure of target construct while offloading is mandatory
>
> Process returned -1 (0xFFFFFFFF) execution time : 0.160 s
> Press ENTER to continue.
>
Perhaps unified shared memory it is not supported?
My pascal gpu (an 1660 GTX super) is, admittedly, a bit old.
However, according to nvidia, my device has unified shared memory.
https://developer.nvidia.com/blog/unified-memory-in-cuda-6/
In CUDA 6, Unified Memory is supported starting with the Kepler GPU architecture (Compute Capability 3.0 or higher)
> /opt/cuda/extras/demo_suite/deviceQuery
> Detected 1 CUDA Capable device(s)
> Device 0: "NVIDIA GeForce GTX 1660 SUPER"
> CUDA Driver Version / Runtime Version 12.4 / 12.6
> CUDA Capability Major/Minor version number: 7.5
</pre>
<img width="1" height="1" alt="" src="http://email.email.llvm.org/o/eJycVl2P27oR_TX0y8CCRK299oMfvPY6XRS53SbZi74FFDmWeEuRKj_sOL--GEr27t4kBXoFAxbEmTnDM4czFCHo1iJu2OKBLfYzkWLn_KYJskvm-6qaNU5dNi_26HxMVkQ0F8Z34Aa0_QA6wAn9BTyG6LWMqEBb0DEAK7f043VE0QdQmgyaFBFCFBF7tLEAHSGIS2D19mrOn44gYPTx2GpnCcNiGEMHrZDWhW8xTgaUTuwQpPMew-Cs0ra9mkhnQ_RJRuhTiGBd_hSFtiDs5TWXQFEUSiO8iNrZAM6D0h5l1CcM4FLM2O74M6yc7g2qYJyP-_nS6QAGhQoQ3Y2kHP7cadmB8BQyos2ZZSoHL2TUUphijPEPiyBMRG8FpUJ5du6MJ_T0enbJKGgwZ5WsPmpUEDrhUUGPvfMXSlkmJRg_iF5lnwytA4iT0EY0BqloU0VPWtxKN3jR9gJcP4DH_yTtMVwxvo4YX0eM0eOCESLtOOh-MAiDd60XPRyFNqGAJ2Lbx4m_ftAGPZx17LITq7cA8xAVq_eS8QfGH3gJMD9OeV1f5mNhA6v39jTEb8u7uT1ppcWcNglz07p-oHyW5fS77kZbaZJCYJy7fii6qUpvFupdL2JXdKx-_GHlhDI6_7MVm3r0Wo5LKZAgrOgxDEKS2BWrH_4So9pG6IW2jK8YX5P__RRojLpl9faa1U651Bhk9SOcOOOrqmR8DSPyzVi7KIDx1e1Dg20OTy5rUsZtBa16-52i3bbxv9En8N2bBFi5nQz4VlWs3p8KJaKY9pVNXtc5rfMfDBg_MH6wDnoxDJliRDV2m6vqX6WYbd-zPbUDtBE9UHAKxPgqOhKeqtjioWT1tirZYs_4TvH3H9Y_lG8K-OfuRvwAwNF5Ilp_x68RNKv3FasfNKt3VZlfsr7HsGQ-VhamJ4Nrwq33ObP8zh-qkYrRYZ8Fsf-J0utHkvizSa22T7Tfo5DIOAf0nqq1hYPQJnmknhQuVnbeWf2dtOpR9KQQq01mfguP5EMsy_Q5L39-daB1YUEbg60w13YjpMQQ4CwCoJUuUQaopsRcP1wrcU1m52xIJkIX45AnQa7eeNgLY0594XzL-EEhDSrGD5-SjbrHUHSxN8Q0KGxS25Iq3JC7a_FruM8ueYlgnMx9HjRNtn58px48eAx5Nu3GFjV2qHlL82DeGm1xHqllhrmz5vIj0FFEYUY4qAjwOJFNk-PPM-ncEYI7Ho0TeZLoAL2wSsRJyVTL8e_Zu8yrx5i8RQXziipVfjtMD503APyGMuXNEEdACZRFtSwh3OJQlMffvjx-ovrTMNQ2YfEGrdw-o-_EEH4xUHTMM9lFCGkYnI-oWH0YPT9eYBBBCgPtkChBYaFaLkv48OVfZE5Taw2g87wVqtcxohqvFAIaHcEZVUxZ_O11zAkpnR9nrYOx39Pn_gIKT1oidOIX2U5T9L26FJ7QuAF9McYqpOsZPzTGkdKmOPMxwFzbPFrmS_Itt08Wdi_7LSwpgZcJ8uPETHilZJx2lPI04xD-jgPNvA_PLyC87HREGUkajK9IbXQ52olBNNroeIG6KEl0nW67zNrb431wQ2T8MM11_Ba9CHlfvfsako44blJL_GeiS8XkuEeCRAXVuIeMZnAikTrWhJNtM7HUBKmf_Pb70_5pCx_w4Oj8UDlzXT-_PD9-mkZp_QgwBt57fUIPv6MPpEXGDzAd29u321Px4i5bVLxYvg_zho2P4g_qqYeP2jpPN6UcxKa-wXyuAeC-WIwczdSmVut6LWa4qe7r9bK-K6vFrNvIJedc1eVdfeTN3bE8HqvVStbrRVkvynXFZ3rDS74oebWoFuRTiBVWslndL9dKCVTI7krshTa3zjTTISTcVPy-ru9mRjRoQr5Lc27xDHmV6FnsZ35DTvMmtYHdlUaHGF7DRB0Nbv6fKwIoh9MxROzpYJyd__csebN5r_ZWxy41k8YJcPqbD979gZKElNMkBU37OG34fwMAAP__XAnlrg">