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

    <tr>
        <th>Summary</th>
        <td>
            HIP attempts no sanity checking on hipLaunchKernel declaration and asserts
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
            clang
      </td>
    </tr>

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

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

<pre>
    Trying to write a minimal reproducer forces you to have a declaration of hipLaunchKernel, but this just asserts if you don't match the signature:

```
// RUN: clang++ -x hip  -nogpuinc -nogpulib -S -emit-llvm --offload-arch=gfx90a %s

void hipLaunchKernel();
__attribute__((global)) void test() { }
```

```
1.      <eof> parser at end of file
2.      /Users/matt/Desktop/hip-assert.hip:5:30: LLVM IR generation of declaration 'test'
3.      /Users/matt/Desktop/hip-assert.hip:5:30: Generating code for declaration 'test'
 #0 0x000000010723fed4 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/usr/local/bin/clang-19+0x102373ed4)
 #1 0x000000010723e3d0 llvm::sys::RunSignalHandlers() (/usr/local/bin/clang-19+0x1023723d0)
 #2 0x000000010724055c SignalHandler(int) (/usr/local/bin/clang-19+0x10237455c)
 #3 0x0000000185c21a24 (/usr/lib/system/libsystem_platform.dylib+0x18046da24)
 #4 0x0000000185bf1cc0 (/usr/lib/system/libsystem_pthread.dylib+0x18043dcc0)
 #5 0x0000000185afda40 (/usr/lib/system/libsystem_c.dylib+0x180349a40)
 #6 0x0000000185afcd30 (/usr/lib/system/libsystem_c.dylib+0x180348d30)
 #7 0x000000010ab15cac (anonymous namespace)::CGNVCUDARuntime::emitDeviceStub(clang::CodeGen::CodeGenFunction&, clang::CodeGen::FunctionArgList&) (.cold.35) (/usr/local/bin/clang-19+0x105c49cac)
 #8 0x00000001076ecb34 (anonymous namespace)::CGNVCUDARuntime::emitDeviceStub(clang::CodeGen::CodeGenFunction&, clang::CodeGen::FunctionArgList&) (/usr/local/bin/clang-19+0x102820b34)
 #9 0x00000001079299ec clang::CodeGen::CodeGenFunction::GenerateCode(clang::GlobalDecl, llvm::Function*, clang::CodeGen::CGFunctionInfo const&) (/usr/local/bin/clang-19+0x102a5d9ec)
#10 0x0000000107949690 clang::CodeGen::CodeGenModule::EmitGlobalFunctionDefinition(clang::GlobalDecl, llvm::GlobalValue*) (/usr/local/bin/clang-19+0x102a7d690)
#11 0x00000001079431a8 clang::CodeGen::CodeGenModule::EmitGlobalDefinition(clang::GlobalDecl, llvm::GlobalValue*) (/usr/local/bin/clang-19+0x102a771a8)
#12 0x0000000107946bbc clang::CodeGen::CodeGenModule::EmitGlobal(clang::GlobalDecl) (/usr/local/bin/clang-19+0x102a7abbc)
```

I'd expect this to just insert the declaration it needs if it doesn't find one already, but it needs to at least error that the signature is not expected 
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzUV81u27wSfRp6Q9igSP0uvEjiOg1uWhRN220wIkcWW4k0SCqN3_6CktNEvmmb5AIf8AVBYsmjOXNmDskj8F7vDOKaZOeE88O9h77TA-GcZJsFDKG1bg3Oo-kXtVWH9Rd30GZHg6U_nQ5Igfba6B466nDvrBokOtpYJ9HTgx1iYAt3MU6h7MBB0NZQ29BW769hMLL9DzqDHeEXtB4CDa329PvgAwXv0QVPdTMmUtYQXgTaQ5AtDS3SWDmEwSERZ4RtCHv4m7Pj73TJt4Rv6eevH4k4o7IDsyP8nPBzuryPVVC6NHa3H7SRx0-drunyhi6x12HZdXc9XS5t03QW1BKcbInY7Jr7igElPPNPoe-sVv_LrCS8IuJ8Crm9hRCcroeAt7fjd-WuszV0MYhXdEwR0IfpOUqKc0qKzfPUnruZrAiriLhA2xDxju7j9ByFQNGo2PhGdzhF8jGSb796dJ7wbQ8hEL7doP8R7J7wbav3y2kMq1bviTjLiDgTLPbx-vrbB3r1me7Q4ONQn844TmukUUxo4v9AuzyimB2VVmEU2B-xKOGCUXbPpp-EFVw0qFIapxnlIs78wU8fPjltwk0A-eOLA4mEl49BDn7eWh8cQk94HjWqTRinEoezHbwjfNtZGae3rbUhfDsKbJlUhJ-z-4RxUQhUaRzur8qSk8pQKPZcZZ8HcxM13r0Ho7qxa-VrwblQbAbO5-ApyzJJZzCEl29gmWaZnAGJJ0BlJnkCPJ1n1DXhW3_wAfvpcvp8u-8gNNb1K3UYYyJEydJcAZ83Mp1B1E0iJXshRGgdgjpBEErKebeyGQI0CtIXIsh5bpFWkM5z5ye5pRJvzV0qMc9dPJ0y1EkmQcbcYKw59Hbw1ECPfj8KvprkdnH58dvF102UXdA9TjfjHrjBOy3xJgw14eW0gU4PWIWXaGYX28HIuCSPy-V30Q9hZ253rePCzY96W0nbqZXIXiW_TKaVhLn8ypnOc5S1SP8VHXjhiis5q8V8OVQzyhWvKpS_xT-tdrx53GcxfjlnejmeURuU40n9uFk9sj37E9uLy4fAK9NYKq15A2XIVIWPU4476XyTr9Iqr9jfKH-wauiOw33X6zBRe6hvg402eqL0kgZM979BN-DYg1cRKlResRmh5ISQSKB8G6F_lkiRQDkjwk-I5HX9VzE-T-T35b-uRKjrJ-J5zkldEV4oivd7lEcnGuxkRrWJvmS0nU99hw7UIKrRpOpAlUU_udRGR7NlkEIXz5nDg7399UCw0ZJ1CD5QdM46GloIc19LtafGhmNBqOhCrYWqRAULXCdFkiVFkuTlol1jXhSySRVmCGnVYCbKhKtS5KlCgVgs9JqzeNazMuE8Z-WKsQwbBWWmGpZnVU1Shj3obhUFsbJut9DeD7iukjTPFh3U2PnjS8LRQo9vCG4d45f1sPMkZZ32wT9mCDp0uH5_9YlCCNjvQ6RDPRgdDlS2KH9EQ2fNqWOetRiMengVWAyuW7ch7Ed_NNr6nQ7tUK-kHQ_IqOXp33Lv7HeU0WCOPKLhHKn8NwAA__9EnOJH">