ye-luo added inline comments.
================ Comment at: openmp/libomptarget/src/omptarget.cpp:233 MapperComponents - .Components[target_data_function == targetDataEnd ? I : E - I - 1]; + .Components[target_data_function == targetDataEnd ? E - I - 1 : I]; MapperArgsBase[I] = C.Base; ---------------- ABataev wrote: > ye-luo wrote: > > ye-luo wrote: > > > ye-luo wrote: > > > > ABataev wrote: > > > > > ABataev wrote: > > > > > > ye-luo wrote: > > > > > > > ye-luo wrote: > > > > > > > > ye-luo wrote: > > > > > > > > > ABataev wrote: > > > > > > > > > > ye-luo wrote: > > > > > > > > > > > ABataev wrote: > > > > > > > > > > > > ABataev wrote: > > > > > > > > > > > > > ye-luo wrote: > > > > > > > > > > > > > > ye-luo wrote: > > > > > > > > > > > > > > > ABataev wrote: > > > > > > > > > > > > > > > > ye-luo wrote: > > > > > > > > > > > > > > > > > ABataev wrote: > > > > > > > > > > > > > > > > > > grokos wrote: > > > > > > > > > > > > > > > > > > > What is the current status of the order > > > > > > > > > > > > > > > > > > > of the arguments clang emits? Is it still > > > > > > > > > > > > > > > > > > > necessary to traverse arguments in > > > > > > > > > > > > > > > > > > > reverse order here? > > > > > > > > > > > > > > > > > > Yes, still required > > > > > > > > > > > > > > > > > Based on the conversation in > > > > > > > > > > > > > > > > > https://reviews.llvm.org/D85216 > > > > > > > > > > > > > > > > > This line of code neither before nor after > > > > > > > > > > > > > > > > > the change plays well. > > > > > > > > > > > > > > > > > Shall we fix the order in targetDataEnd first? > > > > > > > > > > > > > > > > This change is part of this patch and cannot be > > > > > > > > > > > > > > > > committed separately. > > > > > > > > > > > > > > > I mean could you fix that issue as a parent of > > > > > > > > > > > > > > > this patch? > > > > > > > > > > > > > > > This change is part of this patch and cannot be > > > > > > > > > > > > > > > committed separately. > > > > > > > > > > > > > > > > > > > > > > > > > > > > If fixing the reordering is part of this patch, I > > > > > > > > > > > > > > should have seen "target_data_function == > > > > > > > > > > > > > > targetDataEnd ?" branches disappear. > > > > > > > > > > > > > Nope, just with this patch. It reorders the maps and > > > > > > > > > > > > > need to change the cleanup order too. > > > > > > > > > > > > It works just like constructors/destructors: allocate > > > > > > > > > > > > in direct order, deallocate in reversed to correctly > > > > > > > > > > > > handle map order. > > > > > > > > > > > The description says that "present and alloc mappings are > > > > > > > > > > > processed first and then all others." > > > > > > > > > > > Why the order of arguments in targetDataBegin, > > > > > > > > > > > targetDataEnd and targetDataUpdate all get reversed. > > > > > > > > > > Because this is for mappers. Mapper maps are ordered by the > > > > > > > > > > compiler in the direct order (alloc, maps, delete) but when > > > > > > > > > > we need to do exit, we need to release the data in reversed > > > > > > > > > > order (deletes, maps, allocs). > > > > > > > > > I was not making the question clear. My question about > > > > > > > > > "reverse" is not about having a reverse order for > > > > > > > > > targetDataBegin. My question was about "reversing" from the > > > > > > > > > the old code. Your change put the opposite order for > > > > > > > > > targetDataBegin, targetDataEnd and targetDataUpdate cases. > > > > > > > > > I was not making the question clear. My question about > > > > > > > > > "reverse" is not about having a reverse order for > > > > > > > > > targetDataBegin. My question was about "reversing" from the > > > > > > > > > the old code. Your change put the opposite order for > > > > > > > > > targetDataBegin, targetDataEnd and targetDataUpdate cases. > > > > > > > > > > > > > > > > typo correction > > > > > > > > I was not making the question clear. My question about > > > > > > > > "reverse" is not about having a reverse order for > > > > > > > > **targetDataEnd**. My question was about "reversing" from the > > > > > > > > the old code. Your change put the opposite order for > > > > > > > > targetDataBegin, targetDataEnd and targetDataUpdate cases. > > > > > > > My separate question specifically for targetDataEnd is the > > > > > > > following. > > > > > > > > > > > > > > In target(), we call > > > > > > > ``` > > > > > > > targetDataBegin(args) > > > > > > > { // forward order > > > > > > > for (int32_t i = 0; i < arg_num; ++i) { ... } > > > > > > > } > > > > > > > launch_kernels > > > > > > > targetDataEnd(args) > > > > > > > { // reverse order > > > > > > > for (int32_t I = ArgNum - 1; I >= 0; --I) { } > > > > > > > } > > > > > > > ``` > > > > > > > > > > > > > > At a mapper, > > > > > > > ``` > > > > > > > targetDataMapper > > > > > > > { > > > > > > > // generate args_reverse in reverse order for targetDataEnd > > > > > > > targetDataEnd(args_reverse) > > > > > > > } > > > > > > > ``` > > > > > > > Are we actually getting the original forward order due to one > > > > > > > reverse in targetDataMapper and second reverse in targetDataEnd? > > > > > > > Is this the desired behavior? This part confused me. Do I miss > > > > > > > something? Could you explain a bit? > > > > > > Yes, something like this. targetDataEnd reverses the order of > > > > > > mapping arrays. But mapper generator always generates mapping > > > > > > arrays in the direct order (it fills mapping arrays that later > > > > > > processed by the targetDataEnd function). We could fix this by > > > > > > passing extra Boolean flag to the generator function but it means > > > > > > the redesign of the mappers. That's why we have to reverse it in > > > > > > the libomptarget. > > > > > You can check it yourself. Apply the patch, restore the original > > > > > behavior in libomptarget and run libomptarget tests. Mapper related > > > > > tests will crash. > > > > Stick with mapper generator always generating mapping arrays in the > > > > direct order. The targetDataMapper reverse the mapping array and then > > > > passes args_reverse into targetDataEnd. Inside targetDataEnd, mapping > > > > Yes, something like this. targetDataEnd reverses the order of mapping > > > > arrays. But mapper generator always generates mapping arrays in the > > > > direct order (it fills mapping arrays that later processed by the > > > > targetDataEnd function). We could fix this by passing extra Boolean > > > > flag to the generator function but it means the redesign of the > > > > mappers. That's why we have to reverse it in the libomptarget. > > > > > > Stick with mapper generator always generating mapping arrays in the > > > direct order. > > > > > > In the targetDataBegin case, targetDataMapper keep direct order args and > > > calls targetDataBegin(args) and targetDataBegin process args in direct > > > order. > > > > > > In the targetDataEnd case, targetDataMapper reverses the mapping array > > > and then passes args_reverse into targetDataEnd. Inside targetDataEnd, > > > args_reverse are processed in reverse order. So targetDataEnd is actually > > > processing the args in original direct order. This seems contradictory to > > > the constructor/deconstructor like behavior that all the mappings must be > > > processed in the actual reverse order in targetDataEnd. > > > > > > This is my understanding. The current code should be wrong but obviously > > > the current code is working. So why the current code is working? what is > > > inconsistent in my analysis. Could you point out the missing piece. > > > You can check it yourself. Apply the patch, restore the original > > > behavior in libomptarget and run libomptarget tests. Mapper related tests > > > will crash. > > > > For sure without this line, tests would crash and that is why you included > > this line of change in the patch. Since you made the change, you could > > explain why, right? > I changed and simplified codegen for the mapper generator without changing > its interface. I could do this because of the new ordering, before we could > not rely on it. But it also requires a change in the runtime. > targetDataEnd calls targetDataMapper and targetDataMapper fills the array in > the direct order, but targetDataEnd processes them in the reverse order, but > mapper generator does not know about it. It also has to generate the data in > the reverse order, just like targetDataEnd does. > > Before this patch mapper generator tried to do some ordering but it was not > always correct. It was not expecting something like map(alloc:s) map(s.a) > because it was not allowed by the compiler. That's why it worked before and > won't work with this patch. > PS. The change in the mapper generator is also required and cannot be > separated. Without this mappers tests won't work. I played a bit with your patch. ``` #pragma omp target exit data map(from: c.a[0:NUM], c.b[0:NU2M]) map(delete: c) ``` I put NUM=1024 and NU2M = 2048. LIBOMPTARGET_DEBUG reports ``` Libomptarget --> Entry 0: Base=0x00007fff064080a8, Begin=0x00007fff064080a8, Size=16, Type=0x0, Name=(null) Libomptarget --> Entry 1: Base=0x00007fff064080a8, Begin=0x0000000000f9cbd0, Size=4096, Type=0x1000000000012, Name=(null) Libomptarget --> Entry 2: Base=0x00007fff064080b0, Begin=0x0000000000f86e10, Size=8192, Type=0x1000000000012, Name=(null) Libomptarget --> Entry 3: Base=0x00007fff064080a8, Begin=0x00007fff064080a8, Size=16, Type=0x1000000000008, Name=(null) ``` Since targetDataEnd internally reverse the processing order, could you confirm that the frontend was emitting entries 3,2,1,0? I'm wondering if the frontend could emit 3, 0, 1, 2 so the processing order is 2,1,0,3? The spec requires struct element processed before the struct in "target exit data" Repository: rG LLVM Github Monorepo CHANGES SINCE LAST ACTION https://reviews.llvm.org/D86119/new/ https://reviews.llvm.org/D86119 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits