<table border="1" cellspacing="0" cellpadding="8">
    <tr>
        <th>Issue</th>
        <td>
            <a href=https://github.com/llvm/llvm-project/issues/66687>66687</a>
        </td>
    </tr>

    <tr>
        <th>Summary</th>
        <td>
            OpenMP LTO on CUDA target fails with "inconsistent kernel function annotation" assertion
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
            new issue
      </td>
    </tr>

    <tr>
      <th>Assignees</th>
      <td>
      </td>
    </tr>

    <tr>
      <th>Reporter</th>
      <td>
          markdewing
      </td>
    </tr>
</table>

<pre>
    Test case consists of a CUDA file that calls thrust::sort and a main file that uses OpenMP offload to move data to  the device and call sort in the first file.

When compiled with -foffload-lto, the link stage fails with the assertion:

`clang-linker-wrapper: /home/mdewing/prev/software/llvm/github/llvm-project/llvm/lib/Transforms/IPO/OpenMPOpt.cpp:5829: llvm::omp::KernelSet llvm::omp::getDeviceKernels(llvm::Module&): Assertion 'isKernel(*KernelFn) && "Inconsistent kernel function annotation"' failed.`

When compiled without -foffload-lto, it compiles and runs.

It affects clang 17 and mainline.

[lto_error.tar.gz](https://github.com/llvm/llvm-project/files/12652386/lto_error.tar.gz)

</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJyUVNGOqzYQ_RrzMgqCIRB44CH3Rkir9mqv1K36WDkwgLvGRvawUfv1lU1uN9tWVe9LFGvOnMmccybSezUZolaUn0R5SeTGs3XtIt3rQDdlpuRqh9_bF_IMvfQEvTVeefZgR5Dw-efLGUalCXiWAaG1B57d5lkUZ1GcvXUM0gwgYZHKPGA3Tx6eVzJfvoIdR23lAGxhsW8Eg2QZHsAzwUBvqqdIEvghUioTa6NyniNpKrKLyM775y8zGejtsipNA9wUz3AY70MOmq3Az7FdK_MKnuVEMEql_Q4NFek9OVbWhDUemEWV9Vqa6RBayR1uTq4rOVGcQWA324UEdssuncBudfQmsPN25Jt0oab12yKwmxTP2_X-PqzO_kY9v5e1CrUXJ40frVu8wO7p67PAbhfseeW0X1dRnMsamzA89kXF7bLuX34gZ0j_RPxvxYn4EnXdUV5g_Y76YodNk8BK7OTnb2KAwJPye4vAWuB9SGcENhAbKhCIT-YeEzIMrxEC42b6yCGNsSyjtogCT1F6GlJRZf_tod34nzYq_gbyMSJuM_5DFp4Y5DhSzx6ic5CfIjDEUSvzMTii_KTZ_krOWZeydOn0hygvAuuZefVBHez-ci_t7fJg2UcjQyaDbTlWJRZ1FRB_Z8Zmn5oMbTE0RSMTavOqKbPqeKqrZG6P9VjkWVn3eV8d8STLPr_WWSaLkeqmqPJEtZhhkTV5nTd5c6zSgermJHuicaCmH1EcM1qk0mn4eal1U6K836itqqo-JVpeSft4-4iGbhCLwZbykrg2rnTdJi-OmQ5H_87CijW19_P98eUZrNn_DFi6ifjxnASi-t95eD-8ZHO6_W7Z4wJB97jgnwEAAP__GFOXiw">