<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/64738>64738</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
[OMPT] Target callbacks use wrong device number due to late initialization
</td>
</tr>
<tr>
<th>Labels</th>
<td>
new issue
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
Thyre
</td>
</tr>
</table>
<pre>
**Note:** The issue was initially discussed in the ROCm-Developer-Tools/aomp repository. You can find the issue [here](https://github.com/ROCm-Developer-Tools/aomp/issues/527).
There's also a Phabricator review by @mhalk already: https://reviews.llvm.org/D157605
**Description**
In the recently upstreamed implementation of the OpenMP target callbacks of the OMPT interface, it was discovered that the device numbers might not get set in some cases. This seems to affect the `ompt_callback_target` and `ompt_callback_target_data_op` callbacks in particular.
This is bad for tool developers, since we require the correct device number to know which regions were executed on which device.
At the same time, the order of the callbacks seems to be messed up as well. We do see the first target event before the device gets initialized.
**Reproducer**
One can use the following test to reproduce the issue:
```c
#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;
}
```
I've used the callback interface from one of the aomp tests to get the callback information. It can be found [here](https://raw.githubusercontent.com/ROCm-Developer-Tools/aomp/aomp-dev/test/smoke/veccopy-ompt-target-emi/callbacks.h).
Running the tool, we see the following results:
```console
$ clang --version
clang version 18.0.0 (https://github.com/llvm/llvm-project.git 5816d2ab287ab9d2e1624852946973ed43a0e3f2)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/software/software/LLVM/git/bin
$ wget https://raw.githubusercontent.com/ROCm-Developer-Tools/aomp/aomp-dev/test/smoke/veccopy-ompt-target-emi/callbacks.h
$ clang -fopenmp -fopenmp-targets=nvptx64 reproducer.c
$ ./a.out
Callback Target EMI: kind=2 endpoint=1 device_num=-1 task_data=0x55ff20a4aa00 (0x0) target_task_data=0x55ff20a74158 (0x0) target_data=0x7fd23d6287a8 (0x8000000000000001) code=0x55ff1f5497f1
Callback Init: device_num=0 type=sm_75 device=0x55ff20a8d120 lookup=0x7fd23d8730d0 doc=(nil)
Callback Load: device_num:0 filename:(null) host_adddr:0x55ff1f54a668 device_addr:(nil) bytes:613024
Callback DataOp EMI: endpoint=1 optype=1 target_task_data=0x55ff20a74158 (0x0) target_data=0x7fd23d6287a8 (0x8000000000000001) host_op_id=0x7fd23d6287c0 (0x8000000000000002) src=0x7ffc7782efa0 src_device_num=1 dest=(nil) dest_device_num=0 bytes=40 code=0x7fd23d77e393
Callback DataOp EMI: endpoint=2 optype=1 target_task_data=0x55ff20a74158 (0x0) target_data=0x7fd23d6287a8 (0x8000000000000001) host_op_id=0x7fd23d6287c0 (0x8000000000000002) src=0x7ffc7782efa0 src_device_num=1 dest=0x7fd206600000 dest_device_num=0 bytes=40 code=0x7fd23d77e393
Callback DataOp EMI: endpoint=1 optype=2 target_task_data=0x55ff20a74158 (0x0) target_data=0x7fd23d6287a8 (0x8000000000000001) host_op_id=0x7fd23d6287c0 (0x8000000000000003) src=0x7ffc7782efa0 src_device_num=1 dest=0x7fd206600000 dest_device_num=0 bytes=40 code=0x7fd23d77e30e
Callback DataOp EMI: endpoint=2 optype=2 target_task_data=0x55ff20a74158 (0x0) target_data=0x7fd23d6287a8 (0x8000000000000001) host_op_id=0x7fd23d6287c0 (0x8000000000000003) src=0x7ffc7782efa0 src_device_num=1 dest=0x7fd206600000 dest_device_num=0 bytes=40 code=0x7fd23d77e30e
Callback Target EMI: kind=2 endpoint=2 device_num=-1 task_data=0x55ff20a4aa00 (0x0) target_task_data=0x55ff20a74158 (0x0) target_data=0x7fd23d6287a8 (0x8000000000000001) code=0x55ff1f5497f1
Callback Target EMI: kind=1 endpoint=1 device_num=0 task_data=0x55ff20a4aa00 (0x0) target_task_data=0x55ff20a74158 (0x0) target_data=0x7fd23d6287a8 (0x8000000000000004) code=0x55ff1f5498d6
Callback Submit EMI: endpoint=1 req_num_teams=0 target_data=0x7fd23d6287a8 (0x8000000000000004) host_op_id=0x7fd23d6287a0 (0x0)
Callback Submit EMI: endpoint=2 req_num_teams=0 target_data=0x7fd23d6287a8 (0x8000000000000004) host_op_id=0x7fd23d6287a0 (0x0)
Callback Target EMI: kind=1 endpoint=2 device_num=0 task_data=0x55ff20a4aa00 (0x0) target_task_data=0x55ff20a74158 (0x0) target_data=0x7fd23d6287a8 (0x8000000000000004) code=0x55ff1f5498d6
Callback Target EMI: kind=3 endpoint=1 device_num=-1 task_data=0x55ff20a4aa00 (0x0) target_task_data=0x55ff20a74158 (0x0) target_data=0x7fd23d6287a8 (0x8000000000000005) code=0x55ff1f549956
Callback DataOp EMI: endpoint=1 optype=3 target_task_data=0x55ff20a74158 (0x0) target_data=0x7fd23d6287a8 (0x8000000000000005) host_op_id=0x7fd23d6287c0 (0x8000000000000006) src=0x7fd206600000 src_device_num=0 dest=0x7ffc7782efa0 dest_device_num=1 bytes=40 code=0x7fd23d787d7f
Callback DataOp EMI: endpoint=2 optype=3 target_task_data=0x55ff20a74158 (0x0) target_data=0x7fd23d6287a8 (0x8000000000000005) host_op_id=0x7fd23d6287c0 (0x8000000000000006) src=0x7fd206600000 src_device_num=0 dest=0x7ffc7782efa0 dest_device_num=1 bytes=40 code=0x7fd23d787d7f
Callback DataOp EMI: endpoint=1 optype=4 target_task_data=0x55ff20a74158 (0x0) target_data=0x7fd23d6287a8 (0x8000000000000005) host_op_id=0x7fd23d6287c0 (0x8000000000000007) src=0x7fd206600000 src_device_num=0 dest=(nil) dest_device_num=-1 bytes=0 code=0x7fd23d77f75a
Callback DataOp EMI: endpoint=2 optype=4 target_task_data=0x55ff20a74158 (0x0) target_data=0x7fd23d6287a8 (0x8000000000000005) host_op_id=0x7fd23d6287c0 (0x8000000000000007) src=0x7fd206600000 src_device_num=0 dest=(nil) dest_device_num=-1 bytes=0 code=0x7fd23d77f75a
Callback Target EMI: kind=3 endpoint=2 device_num=-1 task_data=0x55ff20a4aa00 (0x0) target_task_data=0x55ff20a74158 (0x0) target_data=0x7fd23d6287a8 (0x8000000000000005) code=0x55ff1f549956
Callback Fini: device_num=0
```
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzsWV9v47gR_zT0y8CGRP31gx-S6AwscGkW26BFnwxaHNlsKFIlKTvupy8o2ZLt_OklQO9yxQaLtSzOcGZ-85sZWWTWio1CXJDkliTFhLVuq83icXswOFlrflgQekPozV-0QxLd9F_gcYsgrG0R9syCUMIJJuUBuLBlay1yEArcFuHHw109LXCHUjdopo9aS0vokum6AYONtsJpc5jBP3QLJVNQCcU7xX53ktxu0SBJCkLzrXON7XxYErrcCLdt17NS14Qu3zFD6LLby39PaEbofEaCggQ3_f-P3f40s8Ck1cDg-5atjSiZ0wYM7gTuYX0AEgf1lsknYNIg4wcS3cClP72snUm5q2fabAhdFmGSpUFybq7Hr0BbGtE4oVV_41zkW4-cwRKVkwdoG-sMstqDWjcSa1SOeVXQVSf50KC6_w6OmQ06KJmUa1Y-2WH5_vsjCOXQVKxEQu9AuC5tPll6hwY94sx1whx3okRQbb1GY6EWm60DpR34rS06n1era4SSWbQzeNwKCxaxtuA0sKrCst-IpIGuG7c6ubPq3SNpAEzxt5ZXnDm20o0XGwMRChpmnChbycxV9oQFYWHNOFTagNNa-hh6HlgfrBWqRNh7QP_VCoOdd6U2xnt6Ea6P4EnpPey3otyCwY3QysIeDQI-Y9k65KDVcblXvfDmpg_dshrBibrD2t_QhqM5pWMMa4BtjVBjVzVtA8xblHIGf0fg2gt1apUw1p1yjDtUDtZY6WM8xzg26IZqFP9GPoOX3PuBjdG8LdG8pN6Dwq4KW3s0qqXUe6E24NBb175me-2xSH0FnFtJg_5febIaCVXKliOQ6E7XzWxLol9eW7OOC_3GKqUDbrMtofTcolAOaiYUoTnstOBA6Py4nt32FwDgSwDuSXIbBr6dRLeDkcawTc3At6QTvr5awHMRatYQmjvtC95rk-im34DO4Z0dRrODD1dyyOquBJ0R69ahZziTEmXHYytqPm7h_yptCM19DAJIVEBAotvu8g7C7prQW0JvxRD7SXHEwLsvSFJ0-mJAYBQszr0u3sPnWbhzeCqj61cAGrcz6FqjOqdPmSmu6HLRAgnNduhpyC9qZmxj4E2CVngqK9aDal1XUd7JK8VKm7prmzP45jqWrz3BW9-M3p4yhu1n_aRpLZpSK4fK_aah4z-mHHeELr1bhC5trZ-Q0OUOy1I3h6lvgNMe0SnWgtDlBcdPc-pHq1RXglvs-pvvKnscG8NQowZtK519sx61slriKa0xlJKpDUynOzTWT6Juob95vAVhPgtmAbw_fP3EO35MG6P_iaXzmEGShymnbE3zjK3nnGKY0jhP6DxO51mEPI5YgFFFB6Y89iMiuoHnPF2l8bRVvh-rqRSqfZ5uVHvq-n4GQ605Si_tHyOeT8PTOl9GvBDGLxG61E0Hvq7cnvlRf375669_u-_DIXS5FmoEZ-8p9FW4cJ2ySjeo6ma4OGpaEhVq17jnNB4btZmVo_7MOzPTretv3Z3Ko0cefrn_5kF7EoqTqKCAijdaKEeiIjzOmJVqaxIV0xAcs0_dvCZRETwnSVXRgMWMBR1fgufAt8jjWH9VNovDJH8pO4hlFacRTz19jmJ5cPEXeq1Scxx2DasknmdVeBXeNyU6Wl2EEIA7NF7V1qssOa6d-5fzkAYgtX5qmzOH8iwKeABclyTyDUMJOTB4sPirZvza4k0AlZCoWN0_SeeqlV4Vttq6FeOce8qOgbA0zU8bsH5xMAfrg0NPzTSMAhqfOu1gv2COPTSnhF7kUTfHuMPfITtdaLpZCX6lVAavKvlmANaUR-mqzLKcYsUCf3N1kT_PSOvOc9DdWF0l-QhUEQcjVXo3sgyjefTboaP_f9D12wZp2u3wP8bvjHr0C-IX_RH4Bfgp_v3E7wq_D0wy-meeZK-GF74zqIMvEF38enQ5T4_kH8L7a7uuhXu1d_jf7z6mVffL6RjZZxx5k9fsDJAXVfmmZ_QP8ewDhKB_PkK8H1301Z9Lk9fDmyfpp4Zl9Dt4-9Fmn142-7Me_qLZB-fN_mwqvGz24XvNPs94Vn1qWP7E7xP4nfEv_oL4ZZ_A773fCdMRu1eeM6osYZ-i3k_oRug-0NS_4iPaf2vqQ3hLocTLlw1XL-MmfBHxeTRnE1yE6ZzSII3y-WS7iOc8wDLDHNd5liVxGgZxzOdlkpb5vOR0IhY0oFGQh2mYhfN4Pot5jimGOM95loY5JXGANRNyOImadK_pF2mcRflEsjVK2532Uapwf3yHTylJiolZdC_x1u3GkjiQwrrxPGvihJPdMeHD_fdHkhSnPI4HGq1F2ButNlfnKrxFcBokczieTnTvYietkYuPvVo8P8_rQvpPAAAA__-a2rEy">