As said on D53250 I think this is the right way to fix these tests: We already do the same for -stdlib=platform and -rtlib=platform.

If I read that patch correctly, this will render -fuse-ld with non-absolute paths useless if a toolchain has DefaultLinker != "ld". I don't think that's what we want to do if the user explicitly sets a different linker.

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.

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.

Ping

Out of interest: Is this fixing a particular issue?

In D52733#1252966, @ABataev wrote:

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?

Even if we will be able to reduce register usage for SPMD, it still going to be high for non-SPMD constructs. The optimizer is not able to understand that it is in parallel region or not during the compilation phase

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.

In D52701#1250645, @grokos wrote:

Looks good. I have no idea why this option was left there...

Going through my list of reviews, this patch was reverted because of memory leaks in other changes. However, I don't think we need this anymore because Clang is raising the PTX level as needed for that CUDA version. Can we abandon this flag?

I think this revision can be closed after rC342924?

I think this revision can be closed after rC342924?

In D52655#1249267, @grokos wrote:

Interesting. Are there other cases where alignment is needed?

In D52434#1249399, @gtbercea wrote:

In D52434#1249186, @Hahnfeld wrote:

In D52434#1249102, @gtbercea wrote:

You report a slow down which I am not able to reproduce actually. Do you use any additional clauses not present in your previous post?

No, only dist_schedule(static) which is faster. Tested on a Tesla P100 with today's trunk version:

#pragma omp target teams distribute parallel for (new defaults)	190 - 250 GB/s
adding clauses for old defaults: schedule(static) dist_schedule(static)	30 - 50 GB/s
same directive with only dist_schedule(static) added (fewer registers)	320 - 400 GB/s

Which loop size you're using ? What runtime does nvprof report for these kernels?

In D52434#1249102, @gtbercea wrote:

You report a slow down which I am not able to reproduce actually. Do you use any additional clauses not present in your previous post?

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.

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].

Ping

Ping

Ping

Ping

In D51875#1249092, @ABataev wrote:

In D51875#1249088, @ABataev wrote:

1. It is not how clang works, it is how standard requires.

In D51875#1249019, @ABataev wrote:

No, you're not correct here.

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;
  }
}

This code is for the distribute loop. And here you have conflict without the datasharing scheme. The problem here is that this check /* received last chunk */ is true for all inner loop iterations for inner for directive and *last_ds may come not from the last iteration of for loop, but from some other iterations. To solve this problem, we need to share the same last_ds between all the threads in the team.

In D51687#1229696, @grokos wrote:

I don't have any comments, just waiting for the --libomptarget-nvptx-path flag to be accepted into the clang driver. Then we can have a look at the rest of this stack.

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.

In D52434#1248974, @gtbercea wrote:

One big problem your code has is that the trip count is incredibly small, especially for STREAM and especially on GPUs. You need a much larger loop size otherwise the timings will be dominated by OpenMP setups costs.

Just tested this and got very weird results for register usage:

void func(double *a) {
  #pragma omp target teams distribute parallel for map(a[0:100]) // dist_schedule(static)
  for (int i = 0; i < 100; i++) {
    a[i]++;
  }
}

In D52434#1248032, @gtbercea wrote:

That is the intention. I just took out that part from here to not confuse things since I wanted to have that as a separate patch. :)

Should we also change the default schedule to static, 1? I know that's not really needed for teams distribute parallel for (because the new default dist_schedule only leaves one iteration per thread), but this doesn't happen for target parallel for. Additionally it would make the intent more explicit and LLVM doesn't need to look through divisions needed to implement static without chunk. Just thinking aloud, not sure if that's worth it.

In D51687#1230375, @gtbercea wrote:

Considering your comment in the description about requiring latest Clang perhaps you should revisit this patch: D46842

Add const per review comments.

(for reference: I missed this in D51623)

Ping

Do we really need new entry points for this? I think we could avoid code duplication by letting the compiler generated code pass the correct chunk to __kmpc_for_static_init_??. This could either come from a (single) new query function (__kmpc_nvptx_distribute_default_chunk?) or can be hard coded to threadsPerBlock because this is only relevant for SPMD.

IMO the current logic makes sense for the end user, prioritizing the environment guess over default paths. I think the tests should be fixed using --cuda-path-ignore-env, I remember I added this parameter to fix the failing tests when I wrote the code.

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.

In D51875#1229536, @ABataev wrote:

In D51875#1229517, @Hahnfeld wrote:

In D51875#1229506, @ABataev wrote:

In D51875#1229502, @Hahnfeld wrote:

In D51875#1229496, @ABataev wrote:

In D51875#1229491, @Hahnfeld wrote:

I really, really dislike adding even more global buffers. 4096 * 32 * 56 are another 7MiB that are not usable for applications. What's wrong with using the existing ones?

Can you upload the CodeGen patch for reductions somewhere? I thought we need a global scratchpad buffer that is adressable for all teams?

I really, really dislike an implementation in ibm-devel, the scratchpad solution will never be added to the trunk. The existing ones cannot be reused, as they are allocated only if the full runtime is used.

What's the overhead of initializing it? The whole libomptarget-nvptx is already a pretty much mess, see my thread on openmp-dev.

It is not the runtime issue, it is the problem with the compiler itself. It breaks compatibility with the other outlined regions and, thus, it cannot be committed to trunk.

Can you please describe the problems? Again, maybe posting the patch may help.

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

In D51875#1229506, @ABataev wrote:

In D51875#1229502, @Hahnfeld wrote:

In D51875#1229496, @ABataev wrote:

In D51875#1229491, @Hahnfeld wrote:

I really, really dislike adding even more global buffers. 4096 * 32 * 56 are another 7MiB that are not usable for applications. What's wrong with using the existing ones?

Can you upload the CodeGen patch for reductions somewhere? I thought we need a global scratchpad buffer that is adressable for all teams?

I really, really dislike an implementation in ibm-devel, the scratchpad solution will never be added to the trunk. The existing ones cannot be reused, as they are allocated only if the full runtime is used.

What's the overhead of initializing it? The whole libomptarget-nvptx is already a pretty much mess, see my thread on openmp-dev.

It is not the runtime issue, it is the problem with the compiler itself. It breaks compatibility with the other outlined regions and, thus, it cannot be committed to trunk.

In D51875#1229496, @ABataev wrote:

In D51875#1229491, @Hahnfeld wrote:

I really, really dislike adding even more global buffers. 4096 * 32 * 56 are another 7MiB that are not usable for applications. What's wrong with using the existing ones?

Can you upload the CodeGen patch for reductions somewhere? I thought we need a global scratchpad buffer that is adressable for all teams?

I really, really dislike an implementation in ibm-devel, the scratchpad solution will never be added to the trunk. The existing ones cannot be reused, as they are allocated only if the full runtime is used.

I really, really dislike adding even more global buffers. 4096 * 32 * 56 are another 7MiB that are not usable for applications. What's wrong with using the existing ones?