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

    <tr>
        <th>Summary</th>
        <td>
            Possible incorrect name mangling in CUDA kernel instantiation
        </td>
    </tr>

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

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

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

<pre>
    I've been diagnosing a problem related to a `cudaKernelLaunch` returning "named symbol not found" for certain kernel launches.
The code that produces the problem is quite complicated and I've been unable to isolate it into something simpler yet, but it seems like the name mangling is producing invalid names.

Running this code:
```
KernelCUDA<ExecutionShapes_t, Threads, Blocks, Functions_t, decltype(Parameters)...><<<Blocks, Threads, MaxShmem, stream>>>(Functions, Parameters...);

if (cudaGetLastError() == cudaErrorSymbolNotFound) {
 auto name = typeid(KernelCUDA<ExecutionShapes_t, Threads, Blocks, Functions_t, decltype(Parameters)...>).name();
    std::cout << "Mangled: " << name << std::endl;
    int status = 0;
    auto demangled = abi::__cxa_demangle(name, 0, 0, &status);
    std::cout << "Demangled: " << demangled << std::endl;
}
```

The mangled name ends up being `FvN5tadma5tupleIJZZZZNS_9OptimizerIJNS_4NodeINS_8SequenceIJEEEZ9InputNodeIJRNS_6TaggedILi0ENS_6TensorIlNS_9AllocatorILNS_6MemoryE1EEENS3_IJLi1ELi10EEEES4_EEEERNS6_ILi2ESC_EERNS6_ILi1ESC_EEEEDaDpOT_EUlDpRKSJ_E_JSD_SF_SH_EEENS2_INS3_IJLi1ELi10ELi768EEEEZNS_6gatherILi0ETkNS_9AnyTensorENS7_IfSA_NS3_IJLi30522ELi768EEEES4_EETkNS_9AnyTensorESD_EEDaRKT0_RKT1_EUlT_RSV_S10_RKT2_E_JNS7_IfSA_SR_S4_EEKSU_KSD_EEENS2_ISR_ZNSS_ILi0ETkNS_9AnyTensorENS7_IfSA_NS3_IJLi2ELi768EEEES4_EETkNS_9AnyTensorESH_EEDaSX_S10_EUlS11_S12_S10_S15_E_JS17_KS1C_KSH_EEENS2_ISR_ZNS_13CombineToNodeITkNS_9AnyTensorENS6_ILm3765833763ES17_EETkNS_9AnyTensorENS6_ILm2639559269ES17_EEZNS_plITkNS_9AnyTensorES1I_TkNS_9AnyTensorES1J_EEDaRKS11_SX_EUlS1M_SX_E_EEDaS1M_SX_S10_EUlS1M_SX_S10_RS13_E_JKS1I_KS1J_S17_EEENS2_ISR_ZNS1H_ITkNS_9AnyTensorENS6_ILm2072007968ES17_EETkNS_9AnyTensorES17_ZNS1K_ITkNS_9AnyTensorES1T_TkNS_9AnyTensorES17_EEDaS1M_SX_EUlS1M_SX_E_EEDaS1M_SX_S10_EUlS1M_SX_S10_S1O_E_JKS1T_KS17_S17_EEENS2_INS3_IJXtlNS_15BlockSizeMarkerILl512EEEEELi10ELi1EEEENS6_ILj2048EZNS_10ReduceNodeILi2ELb1ETkNS_9AnyTensorENS6_ILm3627180312ES17_EEZNHS17_4meanILin1ELb1ERS23_EEDcOSY_EUlS1M_SX_E_UlS1M_E_ZNHS24_ILin1ELb1ES25_EEDcS26_EUlS1M_E_QeqsrSY_6deviceLS9_1EEEDaS10_S15_RKT3_RKT4_EUlS1M_SX_RSY_E_EEJKS23_NS7_IfSA_NS3_IJLi1ELi10ELi1EEEES4_EEEEENS2_IS4_Z10OutputNodeIJRNS6_ILm129000195ENS7_IfSA_NS3_IJLi1ELi10ELi30522EEEES4_EEEEEEDaSL_EUlSO_E_JS2Q_EEEEE12CompileGraphILS9_1EEENS_13CompiledGraphERKSQ_RKS1A_RKS1G_RKS1S_RKS1Y_RKS2M_RKS2T_ENKUlTnivE_clILi3EEEDavENKUlTpTnivE_clIJLi2ELi3EEEEDavENKUlTnivE_clILi0EEEDavEUlRS11_DpOT0_E_ZZZZNS2V_ILS9_1EEES2W_S2Y_S30_S32_S34_S36_S38_S3A_ENKS3C_ILi3EEEDavENKS3E_IJLi2ELi3EEEEDavENKS3G_ILi1EEEDavEUlS3H_S3K_E_EEERNS7_IfNS8_ILS9_3EEESR_S4_EERS1C_RSC_RS17_S3T_E`

And `abi::__cxa_demangle` returns status code -2 `mangled_name is not a valid name under the C++ ABI mangling rules`

This happens when instantiating any kernel at this point in the code, but only if it's instantiated with a lambda.

```
__global__ void RandomKernel(auto) {}

...
RandomKernel<<<1,1>>>([]{}); // Boom
```

So a kernel instantiated with a lambda dependent on a complicated mess of templates breaks name mangling at some point.

I'm still working to isolate the problem and will post any additional findings here, any thoughts appreciated.
</pre>
<img width="1" height="1" alt="" src="http://email.email.llvm.org/o/eJy8V09z46gS_zTKhdqUhGxZPvjg2MqMY8fZMc6-mVwoJHVsNgg0AiWT_fSvGsl_kmzqzelVOQrQdPfvBw10C2vlTgNMguFVMJxfiNbtTTO522kp9O1FbsrXySKgo2cgOYAmpRQ7bazUOyJI3ZhcQUUaUMJBSZwhggRJWLSlWEKjQa1Eq4t9kISkAdc2GvUCSrWooCT2tcqNIto48mhaXQaUkkfTkAIaJ6QmT94EUd4G2MsgnG73QApTAnF74dB_2RZgidvDEYy05GcrHc6raiULj0zokpzTaLXIFSBgaQ2CJ9IRqZ0h1lTg9ojTyqpW0JBXcAGdkbx1OMkCVJYo-QTeKzIhldA7hSrS9ph8Rz8LJUs_xYMPwumm1X4N3F5aTySIcThIwv4XTruFm93Pp0E8y35B0TppNNuLGiz3SLb7BkRpsXmlTPHkW9etLnBiP6eEQrnXGgKa_ikaUYGDxgZ0fHl5GcRZEM-630n_zOit-MX2FVTYtq4BUXkV_6Pp0RGKT7bRMB0H8VVHVD6SgKYYCF_ArYR1WdOYJqBpQMckiOdBPCco9cPMB8LauOsuDMYkGKEdIlpnuiXG-chHlgFN_y9LRMeX6LmD3PEihBDrStyzeFqY1pFuFTGkbzEGAGXYOwh67L591ARdqpNBqR2xTrjWepLhSeLZl1B1lr1U5LKzwXnxS_CDMKBph3VGwuMnoEln97fwzw9-3jE49_8ZjWA0fxfE_VE9qPplAF1a0tYkB38LJOH183roRFmJoWtrBYubh4eHhzXj47vayUr-A83iZs34YG1KWKwZTxn8bEEXsLjJsuxhvNB167zwZrNmPNmK3Q7KxUqGme-CtqZZKLQ4VcoUwplmsULRLVSmec2iLMvWLOaLm5WMspWMwizLMjbg-G-zZglfrCTN2IyfulHXzbK5mNd3W57dq3m9WbIbnvEbNufsmrOv3BumfPHO-kqOkhSVkWayE24Pjce7ffIo9WsHOluzEV88sik_GIjDIaUnfY_xvRKbc4S1WW5DvlluI8S25Rv2F2eRH6GI8Wiabbg3s2T3fOl1O8xswx_WjPHfA_Y_QX31oNh3DyK7VyyKOIuo77Jo6JctGvEli2Z8ebZ0HQwexTNT5VLD1vit_ggIt6WKR8kwjeNREmdo7SOOfh5N4vFwOKbJuJ-HPmr1wSyLFvzj2E2_wJ7D947NrW91JLvOkeixt2FRjESXaHaJdjrn50yjr_wzcjQc0TAcjZP0E3I4iiaWH02waPsvREbneH-bBovuehZbZDF6w6ILiO8Oz1s09Lcuk__ArWieMMjVMKIYJIeDgGev5_c3DQep34go3AC-6H6nfWzl0Wc7GSd0FKVhHNHjTn7FxqACoRcrqSOvvWE0RlbFHfvxhmnXzDiq0QE_aTA69AqMJgeFjH-Dn7ZhP3hSwrMsYMXGPMq6xerCeLPcxvgZnDnZoEueZTdLBPHh6ERv1qK_dvqQGPCHKLxr3fkN52lHdByGYTQefjyKR3vdZXFmE3GuPDC_gYx-68YjOjNVLRV8aUS9XxxYHY4dSkovyjZL9o1j4E_994v_Mv_9gV96679bnq2X92qr5XPGC7VYydh7f-6G66Ogvzni7Fx80gr74Xu1wbOGN22Ie-UfCPoXP0Jl9D-c0R-cxSFnMeUsHnAWJ5zFKWfxFOGweMbfAGFxxv8FAIu_dBf8wTWLv3IWL_0W4v2Py71maeccFQ836Abvrg3DPzwU8ZZnx2dwqkt86z57uI-ZsT0kAT69_YOiUv98cv98SuszZUFOWSVpdQmNz0NnAb0K6BWZXi1O6WjTKrBnD7K0ZC_qGrQlL3vQRGrrhHZSOJ_P69dDyi1cl6HWBrMTqb0Ln672ibDR6pXIRyJdQEf2zBCU5EW6PRFEiSovRZ_3nucGnO-UyYXinDwbWZKN0KWpupwuoClmPYcUsMsrwimmZOH0zcRDBhsFdBad56ddIdNr-8yHBPQ6oNfkypjqY57CsGjpeX_Og5RQgy5BI3ci3hQXFVhLzCNxUNVYS1iSNyCe7LviQDhfXHSL2q8L1iQVsU4qRV5M8-Rrg1NRcl7XYAnzgvNqY53fLFGWEvNZocij1KXUO0v20PhNQrnbm3a3d5aIum6g8KwuL8pJXI7jsbiASTSKh2k0SgeDi_0kz0fJKC3TEgZhnkbJOE6HdBwCTfIBDGlyISc0pMMwCmkUxSM6uAwHj0WY0yJ_FGMQjzQYhFAJqS6Veq4uTbO7kNa2MIloNEjSCyVyUNYXmpRqeCFeGlCKdWczQaU_8nZng0GopHX2ZMZJp2Dyp7FWYtUmdWGaBgr3vvrSBGuCj5spjb5oGzXZO1dbPIY-HnbS7dv8sjBVQK_RV__vj7oxf0PhAnrtEdqAXvcUnif0vwEAAP__pOPT_A">