Paths
| Differential D144569
[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
Comment Actions 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
Revision Contents
Diff 505157 clang/lib/Sema/SemaOpenMP.cpp
clang/test/OpenMP/amdgpu_target_with_aligned_attribute.c
clang/test/OpenMP/parallel_firstprivate_codegen.cpp
clang/test/OpenMP/parallel_master_taskloop_firstprivate_codegen.cpp
clang/test/OpenMP/parallel_master_taskloop_simd_firstprivate_codegen.cpp
clang/test/OpenMP/target_firstprivate_codegen.cpp
clang/test/OpenMP/target_is_device_ptr_codegen.cpp
clang/test/OpenMP/teams_firstprivate_codegen.cpp
| ||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||
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.