ABataev added a comment. In D65835#1623939 <https://reviews.llvm.org/D65835#1623939>, @hfinkel wrote:
> In D65835#1623903 <https://reviews.llvm.org/D65835#1623903>, @ABataev wrote: > > > In D65835#1623883 <https://reviews.llvm.org/D65835#1623883>, @hfinkel wrote: > > > > > In D65835#1623814 <https://reviews.llvm.org/D65835#1623814>, @hfinkel > > > wrote: > > > > > > > In D65835#1623772 <https://reviews.llvm.org/D65835#1623772>, @ABataev > > > > wrote: > > > > > > > > > In D65835#1623756 <https://reviews.llvm.org/D65835#1623756>, @kkwli0 > > > > > wrote: > > > > > > > > > > > In D65835#1622042 <https://reviews.llvm.org/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. > > > > > > I think, this combination is still allowed. I assume that > > > > #Pragma omp target parallel is_device_ptr(a) <dsa_clause>(a) > > > > > > is the same as > > > > #pragma omp target is_device_ptr(a) > > #pragma omp parallel <dsa_clause>(a) > > > > > > i.e. datasharing clauses are applied to inner sub-regions, not the target > > region itself. > > > With the parallel, that makes sense to me. In that case, however, I'd imagine > that the privitization works as normal and the code wouldn't crash. Right? It leads to the use of uninitialized memory. We create private non-initialized copy of the pointer, not the copy of the data it addresses. It shall work with firstprivate clause. Repository: rG LLVM Github Monorepo CHANGES SINCE LAST ACTION https://reviews.llvm.org/D65835/new/ https://reviews.llvm.org/D65835 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
