In D53141#1262378, @Hahnfeld wrote:

In D53141#1262019, @ABataev wrote:

In D53141#1261996, @Hahnfeld wrote:

I guess this will break the case when the DataSize passed to __kmpc_data_sharing_push_stack() needs additional alignment: With this change it is handled in data_sharing_push_stack_common() but __kmpc_data_sharing_push_stack() will determine PushSize without the adjustment and do the final pointer arithmetic.

Why DataSize might require an additional alignment? The DataSize must be already aligned.

It's required in __kmpc_data_sharing_push_stack() for some cases, please see D52655.

Ping!

Ping!

Ping!

Ping!

In D53141#1261996, @Hahnfeld wrote:

I guess this will break the case when the DataSize passed to __kmpc_data_sharing_push_stack() needs additional alignment: With this change it is handled in data_sharing_push_stack_common() but __kmpc_data_sharing_push_stack() will determine PushSize without the adjustment and do the final pointer arithmetic.

LG

In D52733#1252963, @Hahnfeld wrote:

In D52733#1251421, @ABataev wrote:

It might lead to increased register pressure, isn't it? Currently, I'm trying to emit the code that can be optimized out and, thus, may decrease the register pressure. That's why I tried to reduce the number of the runtime checks.

You are right, it's increasing register usage but I think it shouldn't: The generated code is always checking __kmpc_is_spmd_exec_mode first. So if LLVM would be able to optimize this out in SPMD mode, __kmpc_parallel_level should never be called.

I guess this doesn't work because it's illegal to hoist the load of execution_param across a barrier?

LG

It might lead to increased register pressure, isn't it? Currently, I'm trying to emit the code that can be optimized out and, thus, may decrease the register pressure. That's why I tried to reduce the number of the runtime checks.

LG

LG

LG

Why you don't want to use pair<Instruction*, Opcode> as the key in all maps/sets? I expect that it will lead to much more simpler slution.

Extended the test in r343344

In D51875#1249164, @Hahnfeld wrote:

In D51875#1249162, @ABataev wrote:

In D51875#1249159, @Hahnfeld wrote:

In D51875#1249153, @ABataev wrote:

Say, last distribute chunk is [L, U]. In the inner for directive it is split into [L,U1], [U1+1, U2], ..., [Un-1 + 1, U]. Distribute marks all these chunks as last, not the last [Un-1 + 1, U].

I got that. This is why the outer distribute only passes the global address for its last chunk. Then the inner for decides which thread executes [Un-1 + 1, U] and writes the lastprivate value.

Yes, that's right! You got it.

So now you are agreeing to "my" solution which is different than what Clang currently does - I'm confused.

In D51875#1249159, @Hahnfeld wrote:

In D51875#1249153, @ABataev wrote:

In D51875#1249136, @ABataev wrote:

In D51875#1249122, @Hahnfeld wrote:

In D51875#1249092, @ABataev wrote:

In D51875#1249088, @ABataev wrote:

I don't see why the distribute loop cares which thread actually executes the last iteration of the for loop, that's only relevant in the outlined parallel region.

Because it marks as lastprivate not the last loop chunk executed by the last thread, but the set of loop chunks executed by the last team. It means that when you try to write the lastprivate value after the distribute loop you will have multiple writes from the different threads with the different values of lastprivates.

Say, last distribute chunk is [L, U]. In the inner for directive it is split into [L,U1], [U1+1, U2], ..., [Un-1 + 1, U]. Distribute marks all these chunks as last, not the last [Un-1 + 1, U].

I got that. This is why the outer distribute only passes the global address for its last chunk. Then the inner for decides which thread executes [Un-1 + 1, U] and writes the lastprivate value.

Plus, I need to add that I tried the solution you proposed here maybe a month or two ago. If it would work, I would definitely use this one rather than the one implemented now. Because it is much easier to implement and works much faster. But it just does not work!

LG

In D51875#1249136, @ABataev wrote:

In D51875#1249122, @Hahnfeld wrote:

In D51875#1249092, @ABataev wrote:

In D51875#1249088, @ABataev wrote:

I don't see why the distribute loop cares which thread actually executes the last iteration of the for loop, that's only relevant in the outlined parallel region.

Because it marks as lastprivate not the last loop chunk executed by the last thread, but the set of loop chunks executed by the last team. It means that when you try to write the lastprivate value after the distribute loop you will have multiple writes from the different threads with the different values of lastprivates.

In D51875#1249122, @Hahnfeld wrote:

In D51875#1249092, @ABataev wrote:

In D51875#1249088, @ABataev wrote:

I don't see why the distribute loop cares which thread actually executes the last iteration of the for loop, that's only relevant in the outlined parallel region.

In D51875#1249088, @ABataev wrote:

In D51875#1249034, @Hahnfeld wrote:

In D51875#1249019, @ABataev wrote:

Yes, that's the current solution in Clang and actually what I described above:

In D51875#1248997, @Hahnfeld wrote:

In SPMD constructs all CUDA threads are executing the distribute loop, but only the thread executing the last iteration of the for loop has seen the lastprivate value. However the information of which thread this is has been lost at the end of the parallel region. So data sharing is used to communicate the lastprivate value to all threads in the team that is executing the last distribute chunk.

I'm assuming that the pointer returned by /* get data sharing frame from runtime */ is shared between all threads in a team.

1. It is not how clang works, it is how standard requires.
2. Yes, it is shared between all the threads in the team and this is how it is intended to be according to the standard

In D51875#1249034, @Hahnfeld wrote:

In D51875#1249019, @ABataev wrote:

Yes, that's the current solution in Clang and actually what I described above:

In D51875#1248997, @Hahnfeld wrote:

In SPMD constructs all CUDA threads are executing the distribute loop, but only the thread executing the last iteration of the for loop has seen the lastprivate value. However the information of which thread this is has been lost at the end of the parallel region. So data sharing is used to communicate the lastprivate value to all threads in the team that is executing the last distribute chunk.

I'm assuming that the pointer returned by /* get data sharing frame from runtime */ is shared between all threads in a team.

LG

In D51875#1248997, @Hahnfeld wrote:

In D51875#1229536, @ABataev wrote:

I already described it - it breaks the compatibility with other outlined regions and breaks the whole design of the OpenMP implementation.

[...]

Yes, Alex Eichenberger tries to invent something, that will allow us to use something similar to ibm-devel but without breaking the design of OpenMP in the compiler. But it requires some time. But I'd like to have something working, at least.

Just to make sure I came to the right conclusions after trying to understand the code generated since rC342738 and for documentation purposes if the following explanation is correct: The compiler generated code asks the runtime for two loop schedules, one for distribute and the other to implement for. The latter iterates in the chunk returned from the distribute schedule.

For lastprivates on teams distribute parallel for this means that the global value needs to be updated in the last iteration of the last distribute chunk. However, the outlined parallel region only knows whether the current thread is executing the last iteration of the for worksharing construct. This means the lastprivate value of the parallel for is passed back to the distribute loop which decides if it has just executed the last chunk and needs to write to the global value.
In SPMD constructs all CUDA threads are executing the distribute loop, but only the thread executing the last iteration of the for loop has seen the lastprivate value. However the information of which thread this is has been lost at the end of the parallel region. So data sharing is used to communicate the lastprivate value to all threads in the team that is executing the last distribute chunk.

Assume a simple case like this:

int last;
#pragma omp target teams distribute parallel for map(from: last) lastprivate(last)
for (int i = 0; i < 10000; i++) {
  last = i;
}

Clang conceptually generates the following:

void outlined_target_fn(int *last) {
  int *last_ds = /* get data sharing frame from runtime */
  for (/* distribute loop from 0 to 9999 */) {
    outlined_parallel_fn(lb, ub, last_ds);
  }
  if (/* received last chunk */) {
    *last = *last_ds;
  }
}

void outlined_parallel_fn(int lb, int ub, int *last) {
  int last_privatized;
  for (/* for loop from lb to ub */) {
    last_privatized = i;
  }
  if (/* executed last iteration of for loop */) {
    *last = last_privatized;
  }
}

I tried to solve this problem without support from the runtime and this appears to work:

void outlined_target_fn(int *last) {
  int last_dummy;
  for (/* distribute loop from 0 to 9999 */) {
    int *last_p = &last_dummy;
    if (/* is last chunk */) {
      last_p = last;
    }
    outlined_parallel_fn(lb, ub, last_p);
  }
}

void outlined_parallel_fn(int lb, int ub, int *last) {
  int last_privatized;
  for (/* for loop from lb to ub */) {
    last_privatized = i;
  }
  if (/* executed last iteration of for loop */) {
    *last = last_privatized;
  }
}

(Alternatively it should also be possible to set last_p before entering the distribute loop. This will write to last multiple times but the final value should stay in memory after the kernel.)

As you can see the outlined parallel function is unchanged (which is probably what you mean with "breaks the compatibility", @ABataev?). This should work because all invocations of outlined_parallel_fn write their value of last into a dummy location, except the one executing the last distribute chunk.
What do you think?

In D51875#1241913, @grokos wrote:

@Hahnfeld: Are the latest changes in line with your requirements/plans to reduce the memory footprint of the nvptx runtime?

I still think it's a waste of resources to statically allocate around 1 GB on sm_70 / 660 MB on sm_60. And I think it's worrying that we are adding more and more data structures because it seems convenient to quickly solve a problem. The truth seems to be that it's incredibly hard to get rid of them later on...

LG

I forgot to mention that you need an ast print test (the positive test)

LG

LG

In D51937#1247422, @Hahnfeld wrote:

(for reference: I missed this in D51623)

@ABataev can you extend test/mapping/pr38704.c with a map that would have triggered the infinite loop?

LG

LG

Ping!

Ping!

LG

It is not very easy to do, this rework requires redesign. I'm not interested in redesigning the whole library, I have a lot of other work to do.

Fixed message.

Reused preallocated memory for the full runtime as the global memory buffer for the lightweight runtime.

In D51875#1229582, @Hahnfeld wrote:

In D51875#1229565, @ABataev wrote:

In D51875#1229545, @Hahnfeld wrote:

In D51875#1229536, @ABataev wrote:

I already described it - it breaks the compatibility with other outlined regions and breaks the whole design of the OpenMP implementation.

First that's a general statement without any explanation. Second I'm not asking about the scratchpad pointer solution in ibm-devel but rather why we can't pass RequiresDataSharing = true to __kmpc_spmd_kernel_init. Which will give us the data sharing in existing buffers.

First, stop talking like this. I don't owe you anything.

Sorry, my last comment sounds rude even though I didn't mean it.
My point is that it's impossible to review patches without a big picture: what are the other parts, which alternatives did you evaluate, why don't they work?
And to be honest: Disregarding technical review and simply ignoring my comments doesn't feel nice either.