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

    <tr>
        <th>Summary</th>
        <td>
            Crash in Clang generating code for SPIR-V backend
        </td>
    </tr>

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

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

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

<pre>
    I know the SPIR-V backend is experimental, but I figured you might be interested in this crash on 15.0.2.

`saxpy_kernel.cl`:

```opencl
kernel void saxpy(float alpha, global float *x, global float *y,
                  global float *z) {
  uint idx = get_global_id(0);
  z[idx] += alpha * x[idx] + y[idx];
}
```

Command line and crash:

```
$ ./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang -fintegrated-objemitter -target spirv64v1.0 saxpy_kernel.cl -o saxpy_kernel.cl.spv -c -O2                                                                                                  
fatal error: error in backend: cannot select: %24:anyid(s32) = G_FMA %0:id, %21:anyid, %23:anyid (in function: saxpy)
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace, preprocessed source, and associa
ted run script.
Stack dump:
0.      Program arguments: ./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang -fintegrated-objemitter -target spirv64v1.0 saxpy_kernel.cl -o saxpy_kernel.cl.spv -c -O2
1.      <eof> parser at end of file
2.      Code generation
3.      Running pass 'Function Pass Manager' on module 'saxpy_kernel.cl'.
4.      Running pass 'InstructionSelect' on function '@saxpy'
 #0 0x000055c672cd42ef PrintStackTraceSignalHandler(void*) Signals.cpp:0:0
 #1 0x000055c672cd20e4 llvm::sys::CleanupOnSignal(unsigned long) (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x36aa0e4)
 #2 0x000055c672c14432 llvm::CrashRecoveryContext::HandleExit(int) (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x35ec432)
 #3 0x000055c672ccae1e llvm::sys::Process::Exit(int, bool) (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x36a2e1e)
 #4 0x000055c670800116 (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x11d8116)
 #5 0x000055c672c1bcaf llvm::report_fatal_error(llvm::Twine const&, bool) (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x35f3caf)
 #6 0x000055c673a6e65e llvm::reportGISelFailure(llvm::MachineFunction&, llvm::TargetPassConfig const&, llvm::MachineOptimizationRemarkEmitter&, char const*, llvm::StringRef, llvm::MachineInstr const&) (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x444665e)
 #7 0x000055c673a038a1 llvm::InstructionSelect::runOnMachineFunction(llvm::MachineFunction&) (.part.110) InstructionSelect.cpp:0:0
 #8 0x000055c672092b12 llvm::MachineFunctionPass::runOnFunction(llvm::Function&) (.part.65) MachineFunctionPass.cpp:0:0
 #9 0x000055c672511c49 llvm::FPPassManager::runOnFunction(llvm::Function&) (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x2ee9c49)
#10 0x000055c672511ee1 llvm::FPPassManager::runOnModule(llvm::Module&) (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x2ee9ee1)
#11 0x000055c672512987 llvm::legacy::PassManagerImpl::run(llvm::Module&) (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x2eea987)
#12 0x000055c673035d99 clang::EmitBackendOutput(clang::DiagnosticsEngine&, clang::HeaderSearchOptions const&, clang::CodeGenOptions const&, clang::TargetOptions const&, clang::LangOptions const&, llvm::StringRef, llvm::Module*, clang::BackendAction, std::unique_ptr<llvm::raw_pwrite_stream, std::default_delete<llvm::raw_pwrite_stream>>) (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x3a0dd99)
#13 0x000055c673d032ef clang::BackendConsumer::HandleTranslationUnit(clang::ASTContext&) (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x46db2ef)
#14 0x000055c6749c2d01 clang::ParseAST(clang::Sema&, bool, bool) (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x539ad01)
#15 0x000055c673d03df5 clang::CodeGenAction::ExecuteAction() (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x46dbdf5)
#16 0x000055c6736d5251 clang::FrontendAction::Execute() (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x40ad251)
#17 0x000055c67366c67a clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x404467a)
#18 0x000055c6737a4493 clang::ExecuteCompilerInvocation(clang::CompilerInstance*) (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x417c493)
#19 0x000055c670801ce4 cc1_main(llvm::ArrayRef<char const*>, char const*, void*) (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x11d9ce4)
#20 0x000055c6707fa98d ExecuteCC1Tool(llvm::SmallVectorImpl<char const*>&) driver.cpp:0:0
#21 0x000055c6734f90a5 void llvm::function_ref<void()>::callback_fn<clang::driver::CC1Command::Execute(llvm::ArrayRef<llvm::Optional<llvm::StringRef>>, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>*, bool*) const::'lambda'()>(long) Job.cpp:0:0
#22 0x000055c672c142c3 llvm::CrashRecoveryContext::RunSafely(llvm::function_ref<void ()>) (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x35ec2c3)
#23 0x000055c6734fb178 clang::driver::CC1Command::Execute(llvm::ArrayRef<llvm::Optional<llvm::StringRef>>, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>*, bool*) const (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x3ed3178)
#24 0x000055c6734c92d3 clang::driver::Compilation::ExecuteCommand(clang::driver::Command const&, clang::driver::Command const*&, bool) const (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x3ea12d3)
#25 0x000055c6734c9e13 clang::driver::Compilation::ExecuteJobs(clang::driver::JobList const&, llvm::SmallVectorImpl<std::pair<int, clang::driver::Command const*>>&, bool) const (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x3ea1e13)
#26 0x000055c6734d2ddf clang::driver::Driver::ExecuteCompilation(clang::driver::Compilation&, llvm::SmallVectorImpl<std::pair<int, clang::driver::Command const*>>&) (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x3eaaddf)
#27 0x000055c6707ff84e clang_main(int, char**) (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x11d784e)
#28 0x00007fa72e6fb083 __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x24083)
#29 0x000055c6707fa59a _start (./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang+0x11d259a)
clang-15: error: clang frontend command failed with exit code 70 (use -v to see invocation)
clang version 15.0.2 (https://github.com/terralang/llvm-build f67240e51bb7cd80a1f34e948d7a85eb197f59d1)
Target: spirv64v1.0
Thread model: posix
InstalledDir: /scratch2/eslaught/test_spirv/./clang+llvm-15.0.2-x86_64-linux-gnu/bin
clang-15: note: diagnostic msg: 
********************

PLEASE ATTACH THE FOLLOWING FILES TO THE BUG REPORT:
Preprocessed source(s) and associated run script(s) are located at:
clang-15: note: diagnostic msg: /tmp/saxpy_kernel-d29f0e.cl
clang-15: note: diagnostic msg: /tmp/saxpy_kernel-d29f0e.sh
clang-15: note: diagnostic msg: 

********************
```

The requested files in that diagnostic:

[saxpy_kernel-d29f0e.cl.txt](https://github.com/llvm/llvm-project/files/9735894/saxpy_kernel-d29f0e.cl.txt)
[saxpy_kernel-d29f0e.sh.txt](https://github.com/llvm/llvm-project/files/9735896/saxpy_kernel-d29f0e.sh.txt)

Clang version:

```
$ ./clang+llvm-15.0.2-x86_64-linux-gnu/bin/clang --version
clang version 15.0.2 (https://github.com/terralang/llvm-build f67240e51bb7cd80a1f34e948d7a85eb197f59d1)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /scratch2/eslaught/test_spirv/./clang+llvm-15.0.2-x86_64-linux-gnu/bin
```

```
$ uname -a
Linux ... 5.4.0-125-generic #141-Ubuntu SMP Wed Aug 10 13:42:03 UTC 2022 x86_64 x86_64 x86_64 GNU/Linux
```

Build script for Clang:

https://github.com/terralang/llvm-build/blob/f67240e51bb7cd80a1f34e948d7a85eb197f59d1/build.sh

Clang binary at:

https://github.com/terralang/llvm-build/releases/tag/llvm-15.0.2-spirv
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzlWltzo7gS_jXkRWUXd8NDHhznMjmV2aTizO6jS4CwtYOBlUTGmV9_WhLYCOO5xXtqqo7H44AurU9ft7pbiKTK3i7v0eey-oLEhqDl0_3z5E-U4PQzKTNEOSK7mjC6JaXAheUuUNIIdI9yum4YydBb1aAtXW8ESgiipSCMcAHltARx0DtlmG9QVSInmNpTd2rZ15Y9b39Dm-Nd_bb6TFhJimlaQInlzQdt9LeqSQkNVKFuj14rmiElwXKjvKiwQLioN1iiXBdVggukSy13vhsrhH4LLREdfYZtv1pujKzZVde-gckimu2Q5V2jNREr3WFFMwBjQ2PL27f9agVX0NQKrkHSleygcEqxaGfUobf97b6_NbsecNEnaFFttxg0VdCSIHmhGD_FYnvr-mhqubdpgcs1jFoUr9uJVtBkF4Wr0J-AuGY3WZcNNEto2TVGk1wqec0wKHlSJX-TLRWgdDQRmAELiNeUvYb-qzO10UC5aFINi6a8fkWTFE0e3WMNnPuj555jsGNEGKsYkKQvpLW2Fi_LUlyWFUyFFCQVssByA9eHC1y-KfVyz1XWAIq8W91-VA2k4crKhWrtHFrrAq8rgLsIhsubMhW0KqX41oRjDfDp4Wa-vEG8SYBahGG9rREjdcUEEhXaCFFzqV33Fr5rKjZNMk2rLdxILbZ_JjWr_pbg3VvKeUM4XCjjoLCGmoyopa6Xppy3YDglEmkNA7EqJZzDEuZVw3Sx7Ik5r1KKNUa5wllTIp4yWot2US8FiEJZs6331mdPNfVPrAKT2SKwkUZ6EjmD384CNWSnhWx5C1LllneDasw4yAc3ID1ilYPvK4hu7LaNFxVwuiYlAVRSqarSayufm7KkALwGCkH7s9tW9ehJFnzEJV4TBuXSS26rrCmIbDX0jO6spdk_Ifa-5II1SvJSW64W2VmabGP5dmtss9Y3Wa5nI3tnwycI0nDmppnvkhw0Biwrjb5I41jSdYmLD2AIhcQaSdcL3kuuAl3Fp2kt9W6r_3vZzkC2axMfKUMFE_Hm_I3ri0VBcNnUj6WWBiM0JYdLsLOikhYSy3XzaxYDje2dF2IMY-9XmQTnmuAc34eiA7iFXB_PJK1eCXtbVGB0O6FrNA83OyrUYhZngBeQ1Fde5QDPM-GlmDhkjLsnvWL1TR8UxOqqKs7CnQtjG-D8Pjg7sm3HCd85jONkEUgxhgkGKkpSnPc40H5xpXz6Svt0NzpUv3yRQTGtYGVYbnhGQoLcAyAG0rCP1MMhCQNyhPTuHpbmLaYF5E4G0o843QDWzjW0aHszUb5NOgwwRMi9jEkdiXmsBd3Sr8oXPZMtZp9vtJNsO6QbzDoJc1PCUsDCXz-TfFSw8jG9sd9LpO_7IfBkEDkzibS9CDs9JMduTvPblI_lEYvfYVjDB_8upo4jczZ0JH3crUWGWdqxmzjuMVvdWFJtPZSj8E7hCgN5OyJwHFhsAAscJ_XjHrDbJ9m1Czk_D-nXNe0SEgOWvaZlbLCHYAlxfgDsRxUjTeW2RWcCCjgMoM4AqBtHsx7Qgqxx-tY64wPk-21d7GH_i3AxoDHgGnHNs70gi2Oku-gYAc7gSqe7j42oGxkwetXXFK_Ligua8ptyDWbXeY1Dkw8EZ4QtCWbpRjobcAiGR-o1lanRHSm_00r7t-80eoCrsSbf914t3_OBwJaEeWvmC8RFpmuakv7TkFUtwPAWPS-Ov6zqL4wKsgI_QfDW6JSRHDeFWGXgOQT5TkfvRn7fH4qwnYF6Df0biYOX2Z5M6I6nDZGEQzbO-jkNJHslL1Tg-FTSgWHMly9dGnQW5x9miUtyA7qRVvhx6ma204f-JHNxwGECW0KMMwL8uQJ94MUYEBgQgyG7WR6MGPy8296ppIykjSCdnUXn4Q7GNYCZGUiYBeCp-sBumVTe3tz7yM6DycYZDGlgMoN5GMIvNsna1rCbYjLu4hJ2m6N8nZ7EWezQhixkhg3cRoD3Ztj3Y89woRrhAf5rleJjtEfT01umd-J1ZhBJPQNvPEjHnRT2WWnqrLaYmrFnzhh-kx7SW5hpoPRGI6lhb6f37vQ-Tns7MLkBMzIAe5ZDKMtQx-3CeVGLuAd-ucVF8SekZVUbXUfmoCwiYxR2bcc5khzUiOaen8c2DvSzxMNA3aZ5xRRTmoRIPdW70S1SQCKfnazyUsI46FwP3ep_4bTP54bLbVQhh0Id5nBhFB6iWxs8erFntUp3QLG-STCnqQw0VIJSHBmNZcFKMEzlgxhdPRAHk5MGLZ-R7etlk_nBtyqT0MyrLpY7K_A2ybB8uLDnCmba7t__UyXj-jjahrup90Pb8OemXOKcFG8GnyOqQz0859ipA0DDjL2BRSXOLEL_1ybxXpJJBrEjMkj2TZLT2M28kyQrv4uP4lzHvOGkBx3V0_QTGeg3ms4HTxjOQwN2YJYGDcGQBuL8NA2wFPlpDqD2gQL2Eyn2kQveW0iNqTSO9tnTj9LWmtG_Qh5wY5Bnpkh-5mZZfgrode_aCPcjcf4U6f9D8t7t1wjGwIZB12wQnvPIJxpal1d0cJU7mZ8rT5jBOAaQLh2DDGHmkjBP7MhDq1VBE-nTMBMKj3bzt1AKvyMjyeZTXk1DvXP2QYgxSjzMRoIYIy3__XNyQdh-NFUD_a3uBEpe6OONvE1wQcda1zmGDDJDX6jYILKjcl1mBM1sCanhBE1e5bkQJ_L49ZCF9gdCYD-c7s9gZcdvHCMJAIQ1cn2WlDS0ABgQm32bBE6SzNIssrGTez6J_Sib4SggiRPP8iDODum_3tLLifVOZNqqDWyBM3nYQeQTElRXnO50lcqUC5jwNVWkAAieMizSjQuXBPamzXojFEwuVkoy3Pykao50UFZyvz5H2f7RB9rytRq_tY53fHvnsO3Z3vzlZb74gF4-3KDbx4eHx7_u_7hDt_cPN0v08qiKrz7doeebp8fnF6s7S3saO5uLuFxwvfM582Ru34ARpKI31GKxF_mDFADZ21oqonckNcncOLfJtDuSP4MovvkFxZxLPWNH6y8bghj5p9FvNMhTP67fa8CiB2jPZvsbXI3zNBWQtwbX315842e4amj4G8-8IIr9k6pQQ-w92jgQvjkbkPCkIk0g-j2Fvicacna-1xMm3Qi_j_trITelfMmm7EH_fXzhqPmPaKUp8RYiTvsWwIMUhabTKQqm_tSeOG4wUYfgsErVwzxn8ilpStGg5ccn9BesoXmzRo6NHPk2hO_KbaCHPr0skGvDLlADHPy5--MTAFUjfQPqldKR9nkorxhadPlTr9FPK10yVFQymfhx7d-qnntf1rd9YBuzt77__XVcjBQEc7UWBd7XtppWlnBBLp0wDJzQs934Irv0stiL8YWgoiCXak8tfZlG1r25UK51ciEZNF8Eu2hYcfnr754Ekev5F5tLz8dhgNMg8hzseXHo-VGKI9_JM5zlkIBeFDghBb8E32W5bkm-ICUCrsFhXdBLaScOpIEO_LPjqZc7gR27KdzlHk5iC3QEiWAxlTimFVtfsMuWtjWHygK2NfxQCSFTvmJA1HAgHzdiU7FLUhS0EqJdXoRdKAyXag7_BXUDpm4">