[Clang][OpenMP] Clang adding the addrSpace according to DataLayout fix #65483

ajarmusch · 2023-09-06T14:08:59Z

Fix for an issue where clang was not adding the address space according to the data layout, instead was using the default which resulted in a crash at times. The fix includes changes to the cases of LargeCapMemAlloc and CGroupMemAlloc where we are setting the AddrSpace according to the DataLayout.

jdoerfert · 2023-09-06T14:28:07Z

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

@@ -3362,6 +3362,8 @@ Address CGOpenMPRuntimeGPU::getAddressOfLocalVariable(CodeGenFunction &CGF,
      break;
    case OMPAllocateDeclAttr::OMPLargeCapMemAlloc:
    case OMPAllocateDeclAttr::OMPCGroupMemAlloc:
+      if (VD->hasGlobalStorage()) 
+        AS = getLangASFromTargetAS(CGF.CGM.getModule().getDataLayout().getDefaultGlobalsAddressSpace());


This function is only for globals, so we don't need to check VD after all.

I removed the global variable check

…yout fix

alexey-bataev · 2023-09-06T20:45:30Z

Tests?

… DataLayout fix

ajarmusch · 2023-09-11T19:07:35Z

@alexey-bataev sorry for the delay - the test is up

alexey-bataev · 2023-09-11T20:05:33Z

clang/test/OpenMP/test_target_uses_allocators_large_cap_codegen.cpp

+#pragma omp target uses_allocators(omp_large_cap_mem_alloc) allocate(omp_large_cap_mem_alloc: x) firstprivate(x) map(from: device_result)
+{
+  for (int i = 0; i < N; i++) {
+    for (int j = 0; j < N; j++) {
+      x += j + i;
+    }
+  }
+    device_result = x;
+}


Could you simplify the test? We need just an empty target region, everything else can be dropped

Take a look now. Is that reduced enough?

… DataLayout fix

alexey-bataev

LG

jrtc27 · 2023-09-12T19:45:08Z

The title isn't a proper commit subject, the tense/mood doesn't make sense

jrtc27 · 2023-09-12T19:45:50Z

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

@@ -3362,6 +3362,7 @@ Address CGOpenMPRuntimeGPU::getAddressOfLocalVariable(CodeGenFunction &CGF,
      break;
    case OMPAllocateDeclAttr::OMPLargeCapMemAlloc:
    case OMPAllocateDeclAttr::OMPCGroupMemAlloc:
+      AS = getLangASFromTargetAS(CGF.CGM.getModule().getDataLayout().getDefaultGlobalsAddressSpace());


This line is way too long compared to the 80 char line limit

ah I see, what would be the best way to update that line? Push a commit or create another PR?

Running clang-format on your commit (but just your commit, and only accepting changes to the lines you've changed, i.e. just this line) to fix this doesn't need pre-commit review

thanks I updated it

…yout fix (#65483)

jplehr · 2023-09-12T21:30:19Z

I believe this broke one of the AMDGPU OpenMP buildbots https://lab.llvm.org/staging/#/builders/247/builds/6351

dyung · 2023-09-12T22:12:50Z

This change seems to have broken quite a few build bots including:

https://lab.llvm.org/buildbot/#/builders/216/builds/27184
https://lab.llvm.org/buildbot/#/builders/139/builds/49574
https://lab.llvm.org/buildbot/#/builders/247/builds/8920

Could you please take a look and revert if you need time to investigate?

ajarmusch · 2023-09-12T22:21:01Z

Yes I can

…ayout fix (#65483)" This reverts commit e831a32.

gulfemsavrun · 2023-09-12T23:36:11Z

We also started seeing a test failure after this commit:

FAIL: Clang :: OpenMP/test_target_uses_allocators_large_cap_codegen.cpp (13374 of 19273)
--
/b/s/w/ir/x/w/llvm-llvm-project/clang/test/OpenMP/test_target_uses_allocators_large_cap_codegen.cpp:167:18: error: CHECK1-LABEL: expected string not found in input
// CHECK1-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z30test_uses_allocators_large_capv_l102(
                 ^
<stdin>:20:75: note: scanning from here
define dso_local noundef signext i32 @_Z30test_uses_allocators_large_capv() #0 {
                                                                          ^
<stdin>:87:1: note: possible intended match here
define internal void @__omp_offloading_801_2c05cbb__Z30test_uses_allocators_large_capv_l94(i64 noundef %x) #1 {
^

Input file: <stdin>
Check file: /b/s/w/ir/x/w/llvm-llvm-project/clang/test/OpenMP/test_target_uses_allocators_large_cap_codegen.cpp

-dump-input=help explains the following input dump.

Input was:
<<<<<<
             .
             .
             .
            15: @.omp_offloading.entry_name = internal unnamed_addr constant [69 x i8] c"__omp_offloading_801_2c05cbb__Z30test_uses_allocators_large_capv_l94\00" 
            16: @.omp_offloading.entry.__omp_offloading_801_2c05cbb__Z30test_uses_allocators_large_capv_l94 = weak constant %struct.__tgt_offload_entry { ptr @.__omp_offloading_801_2c05cbb__Z30test_uses_allocators_large_capv_l94.region_id, ptr @.omp_offloading.entry_name, i64 0, i32 0, i32 0 }, section "omp_offloading_entries", align 1 
            17: @llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 0, ptr @.omp_offloading.requires_reg, ptr null }] 
            18:  
            19: ; Function Attrs: mustprogress noinline nounwind optnone 
            20: define dso_local noundef signext i32 @_Z30test_uses_allocators_large_capv() #0 { 
label:167'0                                                                               X~~~~~~ error: no match found
            21: entry: 
label:167'0     ~~~~~~~
            22:  %x = alloca i32, align 4 
label:167'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~
            23:  %device_result = alloca i32, align 4 
label:167'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
            24:  %x.casted = alloca i64, align 8 
label:167'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
            25:  %.offload_baseptrs = alloca [2 x ptr], align 8 
label:167'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
             .
             .
             .
            82: omp_offload.cont: ; preds = %omp_offload.failed, %entry 
label:167'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
            83:  ret i32 0 
label:167'0     ~~~~~~~~~~~
            84: } 
label:167'0     ~~
            85:  
label:167'0     ~
            86: ; Function Attrs: noinline norecurse nounwind optnone 
label:167'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
            87: define internal void @__omp_offloading_801_2c05cbb__Z30test_uses_allocators_large_capv_l94(i64 noundef %x) #1 { 
label:167'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
label:167'1     ?                                                                                                                possible intended match
            88: entry: 
label:167'0     ~~~~~~~
            89:  %x.addr = alloca i64, align 8 
label:167'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
            90:  %0 = call i32 @__kmpc_global_thread_num(ptr @1) 
label:167'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
            91:  store i64 %x, ptr %x.addr, align 8 
label:167'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
            92:  %.x..void.addr = call ptr @__kmpc_alloc(i32 %0, i64 4, ptr inttoptr (i64 2 to ptr)) 
label:167'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
             .
             .
             .
>>>>>>

--

https://logs.chromium.org/logs/fuchsia/buildbucket/cr-buildbucket/8770130626803585601/+/u/clang/test/stdout

ajarmusch · 2023-09-12T23:47:25Z

@gulfemsavrun I reverted the commit and am currently trying to fix the issue.

llvm#65483) Fix for an issue where clang was not adding the address space according to the data layout, instead was using the default which resulted in a crash at times. The fix includes changes to the cases of LargeCapMemAlloc and CGroupMemAlloc where we are setting the AddrSpace according to the DataLayout.

…yout fix (llvm#65483)

…ayout fix (llvm#65483)" This reverts commit e831a32.

[Clang][OpenMP] Clang adding the addrSpace according to DataLayout fix

5a56cff

ajarmusch added clang Clang issues not falling into any other category openmp labels Sep 6, 2023

ajarmusch requested a review from jdoerfert September 6, 2023 14:08

ajarmusch requested a review from a team as a code owner September 6, 2023 14:08

ajarmusch self-assigned this Sep 6, 2023

github-actions bot added the clang:codegen label Sep 6, 2023

jdoerfert reviewed Sep 6, 2023

View reviewed changes

fixup! [Clang][OpenMP] Clang adding the addrSpace according to DataLa…

a2403c4

…yout fix

shiltian requested a review from a team September 8, 2023 19:28

fixup! fixup! [Clang][OpenMP] Clang adding the addrSpace according to…

83d92c8

… DataLayout fix

alexey-bataev reviewed Sep 11, 2023

View reviewed changes

fixup! fixup! [Clang][OpenMP] Clang adding the addrSpace according to…

83bdf21

… DataLayout fix

alexey-bataev approved these changes Sep 12, 2023

View reviewed changes

ajarmusch requested review from a team as code owners September 12, 2023 19:31

ajarmusch removed request for a team, PiotrZSL, carlosgalvezp and njames93 September 12, 2023 19:41

ajarmusch merged commit e831a32 into llvm:main Sep 12, 2023
2 checks passed

jrtc27 reviewed Sep 12, 2023

View reviewed changes

ajarmusch added a commit that referenced this pull request Sep 12, 2023

fixup! [Clang][OpenMP] Clang adding the addrSpace according to DataLa…

e3298bb

…yout fix (#65483)

ajarmusch added a commit that referenced this pull request Sep 12, 2023

Revert "[Clang][OpenMP] Clang adding the addrSpace according to DataL…

131ba0a

…ayout fix (#65483)" This reverts commit e831a32.

vzakhari mentioned this pull request Sep 12, 2023

internap proc trampolines #66156

Closed

ZijunZhaoCCK pushed a commit to ZijunZhaoCCK/llvm-project that referenced this pull request Sep 19, 2023

fixup! [Clang][OpenMP] Clang adding the addrSpace according to DataLa…

86ca176

…yout fix (llvm#65483)

ZijunZhaoCCK pushed a commit to ZijunZhaoCCK/llvm-project that referenced this pull request Sep 19, 2023

Revert "[Clang][OpenMP] Clang adding the addrSpace according to DataL…

e32411c

…ayout fix (llvm#65483)" This reverts commit e831a32.

iii-i mentioned this pull request Oct 11, 2023

Add openmp support to System z #66081

Merged

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[Clang][OpenMP] Clang adding the addrSpace according to DataLayout fix #65483

[Clang][OpenMP] Clang adding the addrSpace according to DataLayout fix #65483

ajarmusch commented Sep 6, 2023

jdoerfert Sep 6, 2023

ajarmusch Sep 6, 2023

alexey-bataev commented Sep 6, 2023

ajarmusch commented Sep 11, 2023

alexey-bataev Sep 11, 2023

ajarmusch Sep 12, 2023

alexey-bataev left a comment

jrtc27 commented Sep 12, 2023

jrtc27 Sep 12, 2023

ajarmusch Sep 12, 2023

jrtc27 Sep 12, 2023

ajarmusch Sep 12, 2023

jplehr commented Sep 12, 2023

dyung commented Sep 12, 2023 •

edited

Loading

ajarmusch commented Sep 12, 2023

gulfemsavrun commented Sep 12, 2023

ajarmusch commented Sep 12, 2023

[Clang][OpenMP] Clang adding the addrSpace according to DataLayout fix #65483

[Clang][OpenMP] Clang adding the addrSpace according to DataLayout fix #65483

Conversation

ajarmusch commented Sep 6, 2023

Choose a reason for hiding this comment

Choose a reason for hiding this comment

alexey-bataev commented Sep 6, 2023

ajarmusch commented Sep 11, 2023

Choose a reason for hiding this comment

Choose a reason for hiding this comment

alexey-bataev left a comment

Choose a reason for hiding this comment

jrtc27 commented Sep 12, 2023

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

jplehr commented Sep 12, 2023

dyung commented Sep 12, 2023 • edited Loading

ajarmusch commented Sep 12, 2023

gulfemsavrun commented Sep 12, 2023

ajarmusch commented Sep 12, 2023

dyung commented Sep 12, 2023 •

edited

Loading