This is an archive of the discontinued LLVM Phabricator instance.

[Clang][OpenMP] Fix accessing of aligned arrays in offloaded target regions
ClosedPublic

Authored by doru1004 on Feb 22 2023, 9:13 AM.

Details

Summary

This patch fixes a memory error that occurs when we access an aligned array on the device:

void write_index(int*a, int N) {
    int *aptr __attribute__ ((aligned(64))) = a; // This failed but is fixed by this patch.
    #pragma omp target teams distribute parallel for map(tofrom: aptr[0:N])
    for(int i=0;i<N;i++) {
       aptr[i]=i;
    }
}

Diff Detail

Event Timeline

doru1004 created this revision.Feb 22 2023, 9:13 AM
Herald added a project: Restricted Project. · View Herald TranscriptFeb 22 2023, 9:13 AM
doru1004 requested review of this revision.Feb 22 2023, 9:13 AM
Herald added a project: Restricted Project. · View Herald TranscriptFeb 22 2023, 9:13 AM
jhuber6 added inline comments.Feb 24 2023, 10:13 AM
clang/test/OpenMP/amdgpu_target_with_aligned_attribute.c
2

Can you clang format the test?

jhuber6 added inline comments.Feb 24 2023, 10:15 AM
clang/lib/Sema/SemaOpenMP.cpp
2278

Why does this handling need to be different between CPU and GPU offloading? Strictly speaking, I'm not sure why we need the alignment type here since we'd only get improper alignment on primitive types. So I figured that it should only care about the alignment of the type itself in all cases. Maybe someone can correct me on that.

doru1004 added inline comments.Feb 24 2023, 1:19 PM
clang/lib/Sema/SemaOpenMP.cpp
2278

Are you saying that the previous check was not correct?

jhuber6 added inline comments.Feb 24 2023, 1:27 PM
clang/lib/Sema/SemaOpenMP.cpp
2278

This is the first I've looked at this code, so I don't know what the intention was. But I would assume it's just making sure that the alignment of the uintptr_t is sufficient to contain the by-value copy without causing an addressing error. By that logic I figured it would only care about the alignment of the type, not the declaration itself.

doru1004 added inline comments.Feb 24 2023, 2:11 PM
clang/lib/Sema/SemaOpenMP.cpp
2278

Assuming that what was there before was correct, then you're saying that the Decl type is always the same as Ty. Is that the case?

jhuber6 added inline comments.Feb 24 2023, 2:37 PM
clang/lib/Sema/SemaOpenMP.cpp
2278

I figured that we'd only care about the alignment of the type that's being copied. Because it's not like we can mimic the alignment of the variable on the device. We just need to make sure that its alignment is <= alignof(uintptr_t). Maybe I'm wrong there, someone else could chime in.

doru1004 updated this revision to Diff 505157.Mar 14 2023, 10:20 AM
jhuber6 accepted this revision.Mar 15 2023, 1:24 PM

Pretty sure this is what we discussed in the meeting. So it should be fine to remove the alignment requirement on the declaration since we don't handle it properly by ref either.

This revision is now accepted and ready to land.Mar 15 2023, 1:24 PM
clang/test/OpenMP/parallel_firstprivate_codegen.cpp