OpenMPTarget: scratch implementation for parallel_reduce #3776
Conversation
…s avoided in the OpenMPTarget backend.
|
| while (lock_team != 1) { | ||
| // Avoid tripping over other team's block index. | ||
| // Find a free block index. | ||
| iter = (omp_get_team_num() % max_active_teams); |
There was a problem hiding this comment.
shouldn't this happen outside the while loop? Is that also wrong in the ParallelFor?
There was a problem hiding this comment.
also why do we need two while loops. Isn't the only way to escape the inner while loop for lock_team==1 in which case shmem_block_indx != -1?
There was a problem hiding this comment.
Yes the additional loop is not needed. Deleted it.
| #pragma omp target teams distribute parallel for is_device_ptr(lock_array) | ||
| for (int i = 0; i < max_active_teams; ++i) { | ||
| lock_array[i] = 0; | ||
| } |
There was a problem hiding this comment.
why do this again and again (i probably have overlooked that in the ParallelFor too) I mean we only need to do it when reallocating or (so the get_lock_array could do that). Since otherwise this thing should be back to zero anyway (a kernel will never leave it not being zero unless it crashes).
There was a problem hiding this comment.
Yes, I was initializing it in case a previous kernel crashes. But I guess we don't need since if the kernel crashes, then the application would quit anyway.
| // Loop as long as a shmem_block_index is not found. | ||
| while (shmem_block_index == -1) { | ||
| // Loop until a lock can be acquired on a team | ||
| while (lock_team != 1) { |
…llelReduce` when calculating the shared block index.
Updates in this PR :