No more comments from the community. I think it is okay to accept this revision. Thanks.

I tried the patch in our environment and it works. LG. Thanks.

In D95274#2520002, @JonChesterfield wrote:

What version of cuda are you using? It might be useful to know when things changed.

I think we can dlopen cuda, and try to load cuMemcpyDtoHAsync_v2, and then try to load cuMemcpyDtoHAsync on failure. That would get us some degree of robustness to different symbol versions.

I am seeing these errors.

In D93055#2516351, @vzakhari wrote:

FYI, libomptarget LIT tests fail with profiling enabled in shared-libs builds, because libLLVMSupport.so.12git dependency cannot be resolved. Apparently, LD_LIBRARY_PATH is only set to <build>/llvm/projects/openmp/libomptarget, which obviously does not contain the support library.

It looks fine. Thanks.

default(firstprivate) was added in https://reviews.llvm.org/rG78443666bc18a6957d279a0f58319c8a3e57771a

present map type and motion modifier
default(firstprivate)

In D75001#1922789, @jdoerfert wrote:

In D75001#1891284, @Hahnfeld wrote:

In D75001#1887876, @jdoerfert wrote:

I like this way better. I was hoping we could do it in our cmake only :)

Give others a day or so to comment before your commit but I'm fine with this.

Well, that doesn't really work if openmp-commits is only subscribed on commit. That said, the solution is a bit ugly but I don't have an alternative right now.

What is the problem with openmp-commits here? I got the emails, didn't you?

@kkwli0 Where are we with this?

In D75001#1891966, @tra wrote:

In D75001#1891861, @kkwli0 wrote:

@tra Will it also include -fopenmp-targets=nvptx64-nvidia-cuda?

If you're asking whether the warning will be disabled for OpenMP, then no. This OpenMP target appears to rely on CUDA, and I think all the reasons for the warning existence do still apply.
If you're asking about the workaround, it would only work for non-CUDA compilations that don't really need CUDA SDK, so the answer is probably 'no' as well as you will need a valid CUDA path.

In D75001#1891742, @tra wrote:

In D75001#1891284, @Hahnfeld wrote:

In D75001#1887876, @jdoerfert wrote:

I like this way better. I was hoping we could do it in our cmake only :)

Give others a day or so to comment before your commit but I'm fine with this.

Well, that doesn't really work if openmp-commits is only subscribed on commit. That said, the solution is a bit ugly but I don't have an alternative right now.

Somewhat related, that means Clang issues a warning for every compilation should there be a "unsupported" CUDA version around, even if it's not used? @tra maybe we can only issue the warning if CUDA is going to be used?

This is a good point. It should only affect compilations done with -x cuda. I'll fix that.
Short-term workaround is to specify invalid path to CUDA installation with --cuda-path=/does/not/exist

Somewhat related, that means Clang issues a warning for every compilation should there be a "unsupported" CUDA version around, even if it's not used? @tra maybe we can only issue the warning if CUDA is going to be used?

Since it is a TR8 feature, should we have this guarded by -fopenmp-version=?

So an alternative is to:

• patch openmp/runtime/cmake/LibompCheckLinkerFlag.cmake to make the libomp_check_linker_flag function to ignore the "Unknown CUDA version" warning, AND
• ask users to build with -DCMAKE_CXX_FLAGS=-Wno-unknown-cuda-version -DCMAKE_C_FLAGS=-Wno-unknown-cuda-version to get around the C_SUPPORTS_FPIC test in cmake/modules/HandleLLVMOptions.cmake if the system has CUDA10.2 installed.

Summary

• In order to avoid the bug in PR44587, one needs to build with >9.0.
• With CUDA toolkit 10.2, building the trunk will fail because 10.0.0 does not support it yet.

I am not sure I understand. Do we need to modify stuff outside of /openmp? I was hoping it is our CMake that can be adjusted to make this work as described earlier. TBH, he warning is even not my biggest problem. As long as we get a libomptarget.bc we should be fine.

It turns out that having the warning message also affects the C_SUPPORTS_FPIC test in cmake/modules/HandleLLVMOptions.cmake. As a result, cmake thinks that -fPIC is not supported. Eventually, it leads to error in libclang-cpp.so.

Thanks for all the comments. It makes sense to keep the warning there before any ptx65 features are added. The warning should also apply to both the CUDA and OpenMP compile paths. As a result, I will pursue to ignore the warning in the build of libomp (i.e. libomp_check_linker_flag function in LibompCheckLinkerFlag.cmake). I think it is less pervasive and also can avoid similar occurrence when a new version of CUDA toolkit is available.

In D74571#1874950, @jdoerfert wrote:

in order to avoid bug 44587 which requires >9.0.1 to be the build compiler?

Do we already require CUDA 10 as part of the libomptarget cmake (with a nice message) to avoid the errors you've seen? We for sure need to.

In D74571#1874947, @JonChesterfield wrote:

Do the in tree tests all pass with the 10.2 toolchain? That's not exactly the same as whether it works but is the closest approximation we have available.

Assuming yes, this patch seems uncontroversial.

Change requires unified_address status to partial.

Ping

Address review comments and rebase.

Update based on suggestion to simply the check and rebase.

In D71969#1798224, @jdoerfert wrote:

Is there a reason not to put this check right next to the one that issues err_omp_section_length_negative. SemaExpr.cpp +4668

In D65835#1623883, @hfinkel wrote:

In D65835#1623814, @hfinkel wrote:

In D65835#1623772, @ABataev wrote:

In D65835#1623756, @kkwli0 wrote:

In D65835#1622042, @ABataev wrote:

I want to be sure we're on the same page: For OpenMP 5.0, should we allow is_device_ptr with the private clauses?

Yes, since it is allowed by the standard.

Umm ... I probably missed some earlier discussions! What would be the behavior of the following code?

p = omp_target_alloc(...);
#pragma omp target private(p) is_device_ptr(p)
  p[...] = ...;   // crash or not?

It must crush, I assume. The main problem is that this construct is allowed by the standard.

Yep. We should add a warning message for it.

Upon further reflection, this is not clearly allowed by the standard. My experience is that, when reading standards, sometimes things are disallowed by contradiction (i.e., the standard does not define some behavior, and what the standard does say that's relevant is self contradictory). In this case, 2.19.3 says that list items which are privatized (and which are used) undergo replacement (with new items created as specified) while 2.12.5 says that "The is_device_ptr clause is used to indicate that a list item is a device pointer already in the device data environment and that it should be used directly." A given list item cannot simultaneously be "used directly" (2.12.5) and also undergo replacement: "Inside the construct, all references to the original list item are replaced by references to a new list item received by the task or SIMD lane" (2.19.3). Thus, it may be disallowed.

In D65835#1622042, @ABataev wrote:

I want to be sure we're on the same page: For OpenMP 5.0, should we allow is_device_ptr with the private clauses?

Yes, since it is allowed by the standard.

Looks fine to me.

In D65341#1604562, @kkwli0 wrote:

In D65341#1604561, @gtbercea wrote:

In D65341#1604440, @ABataev wrote:

In D65341#1604411, @Hahnfeld wrote:

There's already D55892 with a better set of tests, including target enter data / target exit data.

Better to merge those two patches into one.

How would you like me to proceed? It looks like the other patch has been sitting there approved for many many months.

Yes, the patch has been waiting for the runtime implementation first and then commit.

In D65341#1604561, @gtbercea wrote:

In D65341#1604440, @ABataev wrote:

In D65341#1604411, @Hahnfeld wrote:

There's already D55892 with a better set of tests, including target enter data / target exit data.

Better to merge those two patches into one.

How would you like me to proceed? It looks like the other patch has been sitting there approved for many many months.

It looks good to me.

It looks good to me.

The change looks okay to me.

I close https://reviews.llvm.org/D57848 and consolidate the review and patch in this one.

Will consolidate the review and patch in https://reviews.llvm.org/D57851 .

See https://reviews.llvm.org/D57848. I think we can close either one.

Committed: r353186

Update based on review comment.

Addressed reviewer's comment.

Add changes in D56790.

In D56733#1360221, @gtbercea wrote:

Could we add the changes in D56790 to this diff?

Add update in OpenMP support release note.