This is an archive of the discontinued LLVM Phabricator instance.

[Clang][OpenMP] Fix failure with team-wide allocated variable
ClosedPublic

Authored by doru1004 on Apr 4 2023, 2:03 PM.

Details

Summary

This patch aims to resolve issue: https://github.com/llvm/llvm-project/issues/60345

The following code:

#include <iostream>
#include <omp.h>
#include <stdlib.h>


int main()
{
	int N =1<<30;
	int *a = new int[N];
#pragma omp target data map(tofrom:a[:N])
	{
           #pragma omp target teams distribute parallel for
            for(int i = 0; i < N; i++)
            {
	       int local_a[10];
               #pragma omp allocate(local_a) allocator(omp_pteam_mem_alloc)
	        for(int j = 0; j < 10; j++)
		    local_a[j] = a[(i+j)%N];
	        a[i] = local_a[0];
  }
	}
std::cout << a[0] << "\n";
}

Fails with the following linker errors:

clang-linker-wrapper: error: <unknown>:0: local_a: unsupported initializer for address space

clang-linker-wrapper: error: Errors encountered inside the LTO pipeline.

Diff Detail

Event Timeline

doru1004 created this revision.Apr 4 2023, 2:03 PM
Herald added a project: Restricted Project. · View Herald TranscriptApr 4 2023, 2:03 PM
doru1004 requested review of this revision.Apr 4 2023, 2:03 PM
Herald added a project: Restricted Project. · View Herald TranscriptApr 4 2023, 2:03 PM
nlopes added inline comments.Apr 4 2023, 2:06 PM
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
3355

Please use poison instead of undef wherever possible as we are tying to remove undef. The replacement is usually safe when you just need a placeholder.
Thank you!

doru1004 updated this revision to Diff 510934.Apr 4 2023, 2:51 PM
Herald added a project: Restricted Project. · View Herald TranscriptApr 4 2023, 2:51 PM
doru1004 marked an inline comment as done.Apr 4 2023, 2:57 PM
doru1004 added inline comments.
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
3355

I've made the change as requested, this also means that I had to add another check in AMDGPUAsmPrinter.cpp.

doru1004 marked an inline comment as done.Apr 4 2023, 2:57 PM
arsenm added inline comments.Apr 4 2023, 3:09 PM
llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
303 ↗(On Diff #510934)

Isa<UndefValue> covers PoisonValue already

doru1004 added inline comments.Apr 4 2023, 3:10 PM
llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
303 ↗(On Diff #510934)

Perfect! I'll revert this part.

doru1004 updated this revision to Diff 510943.Apr 4 2023, 3:16 PM
jhuber6 added inline comments.Apr 5 2023, 8:20 PM
clang/test/OpenMP/target_team_variable_codegen.cpp
34

Shouldn't the Nvidia version also be undefined? Not sure why this should vary depending on the target.

doru1004 added inline comments.Apr 6 2023, 6:37 AM
clang/test/OpenMP/target_team_variable_codegen.cpp
34

Perhaps NVIDIA code path can tolerate a zeroinitializer? I don't want to change it if it's not needed. I am basing this check on the code path for AMD GPUs and the initial bug that was reported.

jdoerfert added inline comments.Apr 6 2023, 8:47 AM
clang/test/OpenMP/target_team_variable_codegen.cpp
34

for AS 3 we should make it always poison.

jhuber6 added inline comments.Apr 6 2023, 8:49 AM
clang/test/OpenMP/target_team_variable_codegen.cpp
34

We should probably change this in HeadToShared in OpenMPOpt as well.

doru1004 added inline comments.Apr 6 2023, 9:08 AM
clang/test/OpenMP/target_team_variable_codegen.cpp
34

Happy to remove the guard and have it always use poison for both NVIDIA and AMD.

doru1004 updated this revision to Diff 511436.Apr 6 2023, 9:14 AM
doru1004 updated this revision to Diff 511444.Apr 6 2023, 9:22 AM
doru1004 marked 3 inline comments as done.Apr 11 2023, 6:27 AM
jhuber6 accepted this revision.Apr 17 2023, 8:05 AM

LGTM unless anyone else has any concerns.

clang/test/OpenMP/target_team_variable_codegen.cpp
34

These should be a single check line now.

This revision is now accepted and ready to land.Apr 17 2023, 8:05 AM