[Regression | OMPT] missing `ompt_scope_end` for `target enter/exit data`

[Thyre]

Description:
The OMPT interface offers different callbacks to track the execution of an OpenMP program. For regions on target devices the callback ompt_callback_target is certainly the most important.

The target region callback is defined like this:

typedef void (*ompt_callback_target_t) (
    ompt_target_t kind,
    ompt_scope_endpoint_t endpoint,
    int device_num,
    ompt_data_t *task_data,
    ompt_id_t target_id,
    const void *codeptr_ra
);

A target region can have different kinds, which describe what type of target directive was used. This issue will focus on two kinds, which were working fine in aomp 16.x and ROCm 5.4.x. We are looking at #pragma omp target enter data and #pragma omp target exit data.

Target region callbacks will get called twice with a different endpoint. When entering a target region, an endpoint of ompt_scope_begin is used. When exiting the region, the endpoint ompt_scope_end is chosen. There's also ompt_scope_beginend but I'm not sure when this is ever used.
Previously, we would get both endpoints.

With this issue present, we are unable to track any target enter data or target exit data regions at all as we require ompt_scope_end to know when the regions did finish! For Score-P, we would run into error messages because we cannot match the locations and wouldn't be able to produce any tracing data.

Reproduce the issue:
We can take the following source code as an example:

#include <omp.h>
#include <stdio.h>
#include "callbacks.h"

int main( void )
{
    int M[10];
#pragma omp target enter data map(to: M[:10]) 
#pragma omp target 
    {
#pragma omp teams distribute parallel for simd
        for(int i = 0; i < 10; ++i)
        {
            M[i] = i;
        }
    }
#pragma omp target exit data map(from: M[:10])
    return 0;
}

callbacks.h is the header file used in the veccopy-ompt-target-tracing example

With ROCm 5.4.x, everything seems fine. However, target_id=6 is missing ompt_scope_end in the device tracing interface:

> ./error_missing_scope_end | grep endpoint
Callback Target: target_id=1 kind=2 endpoint=1 device_num=0 code=0x55da7977084c
Callback Target: target_id=1 kind=2 endpoint=2 device_num=0 code=0x55da7977084c
Callback Target: target_id=4 kind=1 endpoint=1 device_num=0 code=0x55da797708cf
	Record Target: kind=2 endpoint=1 device=0 task_id=0 target_id=1 codeptr=0x55da7977084c
	Record Target: kind=2 endpoint=2 device=0 task_id=0 target_id=1 codeptr=0x55da7977084c
Callback Target: target_id=4 kind=1 endpoint=2 device_num=0 code=0x55da797708cf
Callback Target: target_id=6 kind=3 endpoint=1 device_num=0 code=0x55da7977093a
	Record Target: kind=1 endpoint=1 device=0 task_id=0 target_id=4 codeptr=0x55da797708cf
Callback Target: target_id=6 kind=3 endpoint=2 device_num=0 code=0x55da7977093a
	Record Target: kind=1 endpoint=2 device=0 task_id=0 target_id=4 codeptr=0x55da797708cf
	Record Target: kind=3 endpoint=1 device=0 task_id=0 target_id=6 codeptr=0x55da7977093a

With aomp 17.0, the situation is worse. If we check for the endpoints, all target enter data and target exit data are missing ompt_scope_end

> ./error_missing_scope_end | grep endpoint
Callback Target: target_id=1 kind=2 endpoint=1 device_num=0 code=0x55803e209be0
Callback Target: target_id=4 kind=1 endpoint=1 device_num=0 code=0x55803e209c6c
	Record Target: kind=2 endpoint=1 device=0 task_id=0 target_id=1 codeptr=0x55803e209be0
Callback Target: target_id=4 kind=1 endpoint=2 device_num=0 code=0x55803e209c6c
Callback Target: target_id=6 kind=2 endpoint=1 device_num=0 code=0x55803e209cd8
	Record Target: kind=1 endpoint=1 device=0 task_id=0 target_id=4 codeptr=0x55803e209c6c
	Record Target: kind=1 endpoint=2 device=0 task_id=0 target_id=4 codeptr=0x55803e209c6c
	Record Target: kind=3 endpoint=1 device=0 task_id=0 target_id=6 codeptr=0x55803e209cd8

[Thyre]

The bug is now also present in ROCm 5.6, which is quite unfortunate, since it is a pain / near impossible for us to deal with this bug at configure time. I will check the most recent aomp release..

[Thyre]

Bug is also present in latest aomp release.

[dhruvachak]

Agreed. This needs to be fixed.

[dhruvachak]

This is what we get on trunk with the just-upstreamed implementation.

Callback Target: target_id=0 kind=2 endpoint=1 device_num=-1 code=0x55e4a69245a0
Callback Init: device_num=0 type=gfx90a device=0x55e4a8dad440 lookup=0x7f38ae500c50 doc=(nil)
Callback Load: device_num:0 module_id:0 filename:(null) host_adddr:0x55e4a6920378 device_addr:(nil) bytes:7688
Callback DataOp: target_id=0 host_op_id=1 optype=1 src=0x7ffe0428ee40 src_device_num=1 dest=(nil) dest_device_num=0 bytes=40 code=0x7f38ae3d957f
Callback DataOp: target_id=0 host_op_id=2 optype=2 src=0x7ffe0428ee40 src_device_num=1 dest=0x7f37a1000000 dest_device_num=0 bytes=40 code=0x7f38ae3cfc80
Callback Target: target_id=0 kind=2 endpoint=2 device_num=-1 code=0x55e4a69245a0
Callback Target: target_id=3 kind=1 endpoint=1 device_num=0 code=0x55e4a692462c
Callback Submit: target_id=3 host_op_id=4 req_num_teams=0
Callback Target: target_id=3 kind=1 endpoint=2 device_num=0 code=0x55e4a692462c
Callback Target: target_id=5 kind=3 endpoint=1 device_num=-1 code=0x55e4a6924698
Callback DataOp: target_id=5 host_op_id=6 optype=3 src=0x7f37a1000000 src_device_num=0 dest=0x7ffe0428ee40 dest_device_num=1 bytes=40 code=0x7f38ae3da146
Callback DataOp: target_id=5 host_op_id=7 optype=4 src=0x7f37a1000000 src_device_num=0 dest=(nil) dest_device_num=-1 bytes=0 code=0x7f38ae3de904
Callback Target: target_id=5 kind=3 endpoint=2 device_num=-1 code=0x55e4a6924698
Callback Fini: device_num=0

The target callback endpoints are correct. However, some device numbers are not set.

[mhalk]

First things first: I can reproduce the output from above.
I was very interested in the ordering issue, as Callback Init is printed after(!) Callback Target.

Second, I found the root cause of this in LLVM/main and am further looking into this.
In ASO there's a fix for this very behavior in place, which was commented as temporary.

FYI: when __tgt_target_data_begin_mapper is called the DeviceId is in fact -1.
That's because the CTOR will be called in targetDataMapper but we create the RAII object just before that call.
Unfortunately with the "wrong" information, but there's no other available at this particular point.
The corresponding CTOR call will also trigger OMPT callbacks for Init and Load; then the correct info would be available.

[Thyre]

I've also noticed that the device initialize callback is called right after the first target callback. I don't know if it's still compliant with the OpenMP specification, since it only notes that the initialize callback will be dispatched before any OpenMP construct is executed on the target device.

I've changed how we handle this in Score-P and this shouldn't be an issue for us.

It's good to know that the callbacks are now all dispatched again. Getting the correct device numbers as well would be great since we require that knowledge for proper profiling and tracing, at least during the begin-event.

[mhalk]

This took a bit more effort than anticipated, sorry for the delay.
I used your reproducer and added #pragma omp target update from(M) just before the main's return statement to have an additional example of the target update callback case.

This is the output after applying my fix:

Callback Init: device_num=0 type=gfx90a device=0x992310 lookup=0x7fe1059c0a90 doc=(nil)
Callback Load: device_num:0 module_id:0 filename:(null) host_adddr:0x200378 device_addr:(nil) bytes:7952
Callback Target EMI: kind=2 endpoint=1 device_num=0 task_data=0x4ed480 (0x0) target_task_data=0x50d658 (0x0) target_data=0x7fe1043ed368 (0x8000000000000001) code=0x204a80
  Callback DataOp EMI: endpoint=1 optype=1 target_task_data=0x50d658 (0x0) target_data=0x7fe1043ed368 (0x8000000000000001) host_op_id=0x7fe1043ed380 (0x8000000000000002) src=0x7ffe7e356a30 src_device_num=8 dest=(nil) dest_device_num=0 bytes=40 code=0x7fe1058c974b
  Callback DataOp EMI: endpoint=2 optype=1 target_task_data=0x50d658 (0x0) target_data=0x7fe1043ed368 (0x8000000000000001) host_op_id=0x7fe1043ed380 (0x8000000000000002) src=0x7ffe7e356a30 src_device_num=8 dest=(nil) dest_device_num=0 bytes=40 code=0x7fe1058c974b
  Callback DataOp EMI: endpoint=1 optype=2 target_task_data=0x50d658 (0x0) target_data=0x7fe1043ed368 (0x8000000000000001) host_op_id=0x7fe1043ed380 (0x8000000000000003) src=0x7ffe7e356a30 src_device_num=8 dest=0x7fe0fd000000 dest_device_num=0 bytes=40 code=0x7fe1058b9af4
  Callback DataOp EMI: endpoint=2 optype=2 target_task_data=0x50d658 (0x0) target_data=0x7fe1043ed368 (0x8000000000000001) host_op_id=0x7fe1043ed380 (0x8000000000000003) src=0x7ffe7e356a30 src_device_num=8 dest=0x7fe0fd000000 dest_device_num=0 bytes=40 code=0x7fe1058b9af4
Callback Target EMI: kind=2 endpoint=2 device_num=0 task_data=0x4ed480 (0x0) target_task_data=0x50d658 (0x0) target_data=0x7fe1043ed368 (0x8000000000000001) code=0x204a80
Callback Target EMI: kind=1 endpoint=1 device_num=0 task_data=0x4ed480 (0x0) target_task_data=0x50d658 (0x0) target_data=0x7fe1043ed368 (0x8000000000000004) code=0x204b12
  Callback Submit: target_id=9223372036854775812 (0x8000000000000004) host_op_id=1 req_num_teams=0
Callback Target EMI: kind=1 endpoint=2 device_num=0 task_data=0x4ed480 (0x0) target_task_data=0x50d658 (0x0) target_data=0x7fe1043ed368 (0x8000000000000004) code=0x204b12
Callback Target EMI: kind=3 endpoint=1 device_num=0 task_data=0x4ed480 (0x0) target_task_data=0x50d658 (0x0) target_data=0x7fe1043ed368 (0x8000000000000005) code=0x204b74
  Callback DataOp EMI: endpoint=1 optype=3 target_task_data=0x50d658 (0x0) target_data=0x7fe1043ed368 (0x8000000000000005) host_op_id=0x7fe1043ed380 (0x8000000000000006) src=0x7fe0fd000000 src_device_num=0 dest=0x7ffe7e356a30 dest_device_num=8 bytes=40 code=0x7fe1058cadba
  Callback DataOp EMI: endpoint=2 optype=3 target_task_data=0x50d658 (0x0) target_data=0x7fe1043ed368 (0x8000000000000005) host_op_id=0x7fe1043ed380 (0x8000000000000006) src=0x7fe0fd000000 src_device_num=0 dest=0x7ffe7e356a30 dest_device_num=8 bytes=40 code=0x7fe1058cadba
  Callback DataOp EMI: endpoint=1 optype=4 target_task_data=0x50d658 (0x0) target_data=0x7fe1043ed368 (0x8000000000000005) host_op_id=0x7fe1043ed380 (0x8000000000000007) src=0x7fe0fd000000 src_device_num=0 dest=(nil) dest_device_num=-1 bytes=0 code=0x7fe1058d1fb9
  Callback DataOp EMI: endpoint=2 optype=4 target_task_data=0x50d658 (0x0) target_data=0x7fe1043ed368 (0x8000000000000005) host_op_id=0x7fe1043ed380 (0x8000000000000007) src=0x7fe0fd000000 src_device_num=0 dest=(nil) dest_device_num=-1 bytes=0 code=0x7fe1058d1fb9
Callback Target EMI: kind=3 endpoint=2 device_num=0 task_data=0x4ed480 (0x0) target_task_data=0x50d658 (0x0) target_data=0x7fe1043ed368 (0x8000000000000005) code=0x204b74
Callback Target EMI: kind=4 endpoint=1 device_num=0 task_data=0x4ed480 (0x0) target_task_data=0x50d658 (0x0) target_data=0x7fe1043ed368 (0x8000000000000008) code=0x204bb4
Callback Target EMI: kind=4 endpoint=2 device_num=0 task_data=0x4ed480 (0x0) target_task_data=0x50d658 (0x0) target_data=0x7fe1043ed368 (0x8000000000000008) code=0x204bb4
Callback Fini: device_num=0

Previously kind=2/3/4 did not generate endpoint=2 callbacks, now they are called the same number of times (as expected).

@Thyre Please, if you find the time, check this output and if applicable confirm that this is the desired output.
I will then open a Phabricator review, so this is properly fixed upstream as well as on ASO, and report back here.
Just want to make sure everything's in order before I put it up for review.

Thank you!

[mhalk]

Accidentally put up the mixed EMI (all EMI but Submit, for some testing) output but maybe it is interesting as well, so I will leave it in the previous comment.

Here the non-EMI output:

Callback Init: device_num=0 type=gfx90a device=0x1c95310 lookup=0x7f33d0a90a90 doc=(nil)
Callback Load: device_num:0 module_id:0 filename:(null) host_adddr:0x200378 device_addr:(nil) bytes:7952
Callback Target: target_id=0 kind=2 endpoint=1 device_num=0 code=0x2044d0
  Callback DataOp: target_id=0 host_op_id=1 optype=1 src=0x7ffe350389c0 src_device_num=8 dest=(nil) dest_device_num=0 bytes=40 code=0x7f33d099974b
  Callback DataOp: target_id=0 host_op_id=2 optype=2 src=0x7ffe350389c0 src_device_num=8 dest=0x7f33c3c00000 dest_device_num=0 bytes=40 code=0x7f33d0989af4
Callback Target: target_id=0 kind=2 endpoint=2 device_num=0 code=0x2044d0
Callback Target: target_id=3 kind=1 endpoint=1 device_num=0 code=0x204562
  Callback Submit: target_id=3 host_op_id=4 req_num_teams=0
Callback Target: target_id=3 kind=1 endpoint=2 device_num=0 code=0x204562
Callback Target: target_id=5 kind=3 endpoint=1 device_num=0 code=0x2045c4
  Callback DataOp: target_id=5 host_op_id=6 optype=3 src=0x7f33c3c00000 src_device_num=0 dest=0x7ffe350389c0 dest_device_num=8 bytes=40 code=0x7f33d099adba
  Callback DataOp: target_id=5 host_op_id=7 optype=4 src=0x7f33c3c00000 src_device_num=0 dest=(nil) dest_device_num=-1 bytes=0 code=0x7f33d09a1fb9
Callback Target: target_id=5 kind=3 endpoint=2 device_num=0 code=0x2045c4
Callback Target: target_id=8 kind=4 endpoint=1 device_num=0 code=0x204604
Callback Target: target_id=8 kind=4 endpoint=2 device_num=0 code=0x204604
Callback Fini: device_num=0

[Thyre]

Looks good at first glance. Thanks for looking and fixing the issue.

[mhalk]

Sure & thanks for the quick response!
Phab review opened:
https://reviews.llvm.org/D157605

[mhalk]

@Thyre If it's not too much to ask for, could you please open a corresponding issue for LLVM?
Maybe this would help to bump the priority / visibility -- also, then I could add sth. like fixes issue X.
I've just revisited the Phabricator review, since Dhruva has posted some feedback.

[Thyre]

@Thyre If it's not too much to ask for, could you please open a corresponding issue for LLVM? Maybe this would help to bump the priority / visibility -- also, then I could add sth. like fixes issue X. I've just revisited the Phabricator review, since Dhruva has posted some feedback.

Sure, I'll work on that after my next meeting. I'll also link to this issue to preserve the discussion.

[mhalk]

Thank you very much!