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

    <tr>
        <th>Summary</th>
        <td>
            hipcc crash on __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12
        </td>
    </tr>

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

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

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

<pre>
    
hipcc version 6.3.42134-a9a80e791
os: 24.04.2

```
$ hipcc --version
HIP version: 6.3.42134-a9a80e791
AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.3 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/rocm-6.3.3/lib/llvm/bin
Configuration file: /opt/rocm-6.3.3/lib/llvm/bin/clang++.cfg
```



build command:
```
$ hipcc --offload-arch=gfx1201 --cuda-device-only  -S  wmma_bug.cpp
clang++: warning: argument unused during compilation: '--hip-link' [-Wunused-command-line-argument]
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace, preprocessed source, and associated run script.
Stack dump:
0.      Program arguments: /opt/rocm-6.3.3/lib/llvm/bin/clang++ --offload-arch=gfx1201 -O3 --driver-mode=g++ -O3 --hip-link --cuda-device-only -S -x hip wmma_bug.cpp
1.      <eof> parser at end of file
2.      Code generation
3.      Running pass 'CallGraph Pass Manager' on module 'wmma_bug.cpp'.
4.      Running pass 'AMDGPU DAG->DAG Pattern Instruction Selection' on function '@r_2_2_2_2_2_2_2_2_2_2_2_2'
Stack dump without symbol names (ensure you have llvm-symbolizer in your PATH or set the environment var `LLVM_SYMBOLIZER_PATH` to point to it):
0  clang++ 0x00005fe7cc3851c2 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) + 66
1  clang++   0x00005fe7cc382edc llvm::sys::CleanupOnSignal(unsigned long) + 140
2  clang++   0x00005fe7cc2ba558
3  libc.so.6 0x000072d054245330
4 clang++   0x00005fe7ca9bf218
5  clang++   0x00005fe7cd3e809a llvm::SelectionDAGISel::SelectCodeCommon(llvm::SDNode*, unsigned char const*, unsigned int) + 3034
6  clang++   0x00005fe7cd3e339c llvm::SelectionDAGISel::DoInstructionSelection() + 540
7  clang++ 0x00005fe7cd3ee9ec llvm::SelectionDAGISel::CodeGenAndEmitDAG() + 668
8 clang++   0x00005fe7cd3f12f7 llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) + 4423
9 clang++   0x00005fe7cd3f3886
10 clang++   0x00005fe7cb7bad18
11 clang++ 0x00005fe7cbd25e29 llvm::FPPassManager::runOnFunction(llvm::Function&) + 1097
12 clang++   0x00005fe7cb422441
13 clang++   0x00005fe7cbd26b7e llvm::legacy::PassManagerImpl::run(llvm::Module&) + 1070
14 clang++ 0x00005fe7cc5d3236 clang::EmitBackendOutput(clang::DiagnosticsEngine&, clang::HeaderSearchOptions const&, clang::CodeGenOptions const&, clang::TargetOptions const&, clang::LangOptions const&, llvm::StringRef, llvm::Module*, clang::BackendAction, llvm::IntrusiveRefCntPtr<llvm::vfs::FileSystem>, std::unique_ptr<llvm::raw_pwrite_stream, std::default_delete<llvm::raw_pwrite_stream>>, clang::BackendConsumer*) + 5414
15 clang++   0x00005fe7ccb6613c clang::BackendConsumer::HandleTranslationUnit(clang::ASTContext&) + 1052
16 clang++   0x00005fe7ce36c939 clang::ParseAST(clang::Sema&, bool, bool) + 1177
17 clang++   0x00005fe7cce1cda9 clang::FrontendAction::Execute() + 201
18 clang++   0x00005fe7ccd9da69 clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) + 297
19 clang++   0x00005fe7cceddc9b clang::ExecuteCompilerInvocation(clang::CompilerInstance*) + 523
20 clang++ 0x00005fe7caac148b cc1_main(llvm::ArrayRef<char const*>, char const*, void*) + 4747
21 clang++   0x00005fe7caab9aba
22 clang++   0x00005fe7ccbb51ad
23 clang++   0x00005fe7cc2ba987 llvm::CrashRecoveryContext::RunSafely(llvm::function_ref<void ()>) + 39
24 clang++   0x00005fe7ccbb58d1
25 clang++   0x00005fe7ccb7713a clang::driver::Compilation::ExecuteCommand(clang::driver::Command const&, clang::driver::Command const*&, bool) const + 170
26 clang++ 0x00005fe7ccb77b37
27 clang++   0x00005fe7ccb7d3a1 clang::driver::Compilation::ExecuteJobs(clang::driver::JobList const&, llvm::SmallVectorImpl<std::pair<int, clang::driver::Command const*>>&, bool) const + 10609
28 clang++   0x00005fe7ccb8d9fc clang::driver::Driver::ExecuteCompilation(clang::driver::Compilation&, llvm::SmallVectorImpl<std::pair<int, clang::driver::Command const*>>&) + 364
29 clang++   0x00005fe7caabebbc clang_main(int, char**, llvm::ToolContext const&) + 8316
30 clang++   0x00005fe7ca9e695a main + 106
31 libc.so.6 0x000072d05422a1ca
32 libc.so.6 0x000072d05422a28b __libc_start_main + 139
33 clang++   0x00005fe7caab87d5 _start + 37
clang++: error: clang frontend command failed with exit code 139 (use -v to see invocation)
AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.3 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/rocm-6.3.3/lib/llvm/bin
Configuration file: /opt/rocm-6.3.3/lib/llvm/bin/clang++.cfg
clang++: note: diagnostic msg: 
********************

PLEASE ATTACH THE FOLLOWING FILES TO THE BUG REPORT:
Preprocessed source(s) and associated run script(s) are located at:
clang++: note: diagnostic msg: /tmp/wmma_bug-gfx1201-802f0b.hip
clang++: note: diagnostic msg: /tmp/wmma_bug-gfx1201-802f0b.sh
clang++: note: diagnostic msg: 

********************
failed to execute:/opt/rocm-6.3.3/lib/llvm/bin/clang++  --offload-arch=gfx1201 -O3 --driver-mode=g++ -O3 --hip-link  --cuda-device-only -S -x hip wmma_bug.cpp
```

wmma_bug.cpp
```cpp
typedef long unsigned int size_t;
typedef unsigned short hip_bfloat16;
typedef hip_bfloat16 hip_bfloat168 __attribute__((ext_vector_type(8)));
static inline __attribute__((device)) hip_bfloat168 make_hip_bfloat168(hip_bfloat16 x, hip_bfloat16 y, hip_bfloat16 z, hip_bfloat16 w, hip_bfloat16 a, hip_bfloat16 b, hip_bfloat16 c, hip_bfloat16 d) { return { x, y, z, w, a, b, c, d }; }
#define __WMMA_16_16_16___bf16___bf16 __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12
extern "C" __attribute__((device, const)) size_t __ockl_get_local_id(unsigned int);
extern "C" __attribute__((global))void __attribute__((amdgpu_flat_work_group_size(1, 32)))r_2_2_2_2_2_2_2_2_2_2_2_2(hip_bfloat16* data0, hip_bfloat16* data1, hip_bfloat16* data2) {
  int lidx0 = __ockl_get_local_id(0); /* 8 */
  int lidx1 = __ockl_get_local_id(1); /* 2 */
  int lidx2 = __ockl_get_local_id(2); /* 2 */
  int alu0 = (lidx0>>2);
  int alu1 = (lidx0&1);
  int alu2 = ((lidx0>>1)&1);
  int alu3 = ((lidx2<<2)+(lidx1<<7)+(alu0<<6)+(alu1<<4)+(alu2<<5));
  hip_bfloat16 val0 = *(data1+alu3);
  int alu4 = ((lidx1<<3)+(alu0<<2)+(alu2<<1)+alu1);
  int alu5 = ((lidx2<<6)+alu4);
  hip_bfloat16 val1 = *(data2+alu5);
 hip_bfloat16 val2 = *(data1+(alu3+1));
  hip_bfloat16 val3 = *(data1+(alu3+2));
  hip_bfloat16 val4 = *(data1+(alu3+3));
  hip_bfloat16 val5 = *(data1+(alu3+8));
  hip_bfloat16 val6 = *(data1+(alu3+9));
 hip_bfloat16 val7 = *(data1+(alu3+10));
  hip_bfloat16 val8 = *(data1+(alu3+11));
  hip_bfloat16 val9 = *(data2+(alu5+16));
 hip_bfloat16 val10 = *(data2+(alu5+32));
  hip_bfloat16 val11 = *(data2+(alu5+48));
  hip_bfloat16 val12 = *(data2+(alu5+128));
 hip_bfloat16 val13 = *(data2+(alu5+144));
  hip_bfloat16 val14 = *(data2+(alu5+160));
  hip_bfloat16 val15 = *(data2+(alu5+176));
  unsigned short precast0 = ((unsigned short)(0u));
  hip_bfloat16 cast0 = (*((hip_bfloat16*)&precast0));
  hip_bfloat168 wmma0 = __WMMA_16_16_16___bf16___bf16(make_hip_bfloat168(val0,val2,val3,val4,val5,val6,val7,val8), make_hip_bfloat168(val1,val9,val10,val11,val12,val13,val14,val15), make_hip_bfloat168(cast0,cast0,cast0,cast0,cast0,cast0,cast0,cast0));
  int alu6 = ((lidx2<<7)+alu4);
  *(data0+alu6) = wmma0[0];
  *(data0+(alu6+16)) = wmma0[1];
  *(data0+(alu6+32)) = wmma0[2];
  *(data0+(alu6+48)) = wmma0[3];
  *(data0+(alu6+64)) = wmma0[4];
  *(data0+(alu6+80)) = wmma0[5];
  *(data0+(alu6+96)) = wmma0[6];
  *(data0+(alu6+112)) = wmma0[7];
}
```


</pre>
<img width="1" height="1" alt="" src="http://email.email.llvm.org/o/eJzsOktT47qav0ZsVElZkp8LFk5C6L4FFwo499TMxiVbSuLBkTyyHOD8-ilJTrDzMHSfmVndblc5SN_7pU9fmTZNuRacX4NgBoLFFW31RqrrHOn1VS7ZxzXw0k1ZFwXccdWUUsBwSqY-RsSf0ITGHo8SBLxUNoCkEPtTz59i4KXmCb3u8VKAfejITCYdIeClP34-7ska7POU0_sFLCoq1gcJUDz1pt661BDgeKN1bXgDvAR4uS71ps2nhdwCvHyijEvxUHMxl9u61RzgZVXttpNayf_ihYZKFhPDlEAceAhDHuSriAdBkSA_8XIvKvyExUnorWiUxITFoZ-QsPABToCXvlC15toI_h6HWehPWvEq5JuYVKVo3ydr0RqgjeKUwa1kvDKgtWzKd-ClP0WjaVVxtiiVWQd4KWsN8FLJYuuEMtKWeSczwMu8NDabS7Eq162i2phiVVb8F9Dx0loS4BnAs2mxWh-7af_kbVkxWMjtlgpmrHvZnXK1qiRlE6qKDSCL9eodYQ_ByaRoGZ0wvisLPpGi-oBw8gzh23ZLs7xdT4u6Bl7ak8fo8UaVKMXa_KRq3W650LAVbcMZZK0qxdrIVJeV1d4pHk0mm7I2Rn8FOIIgmE3-dCiTTn6zxyd7eiBYAC99vLtJn29g0-bbUkMK83YNFa-l0lBLOBJUnTX7cQTwsmyaljcALyEVDJaiqFrGod5wWCjabGBOi1etaMEBnsNa8VrJgjdGrUa2yi0bTNo0siip5gyqVsCmUGWtp8BLnzUtXiFrt7XzhjeF9t-jkmtFtwdrNb8bDSOefCBwMmGq3HE1MYFs9vZIdm9v_3NOnzzDybsJlmPPo04DQOZcrgC5gTVVDVeQasgFg3LlottLcQc5l4zDNRfcBT_wUtLtPLXChA2sadOYkJjTqrpVtN7AR7NyTwVdc2XCQwqTi23FDdhAIhwZO_sXKKb3i9vHP-AivZ0AcrNIb-Ej1ZorAU0mq7aw6fjMK25_daxWrXAbAEfA91SGz_83230fw7dSb2SrYfOxzWUFBd1yI0bMRdMqDj9kCzd0x6ENQwdU_sUVLIXZU_AxffkBpYIN1zYKudiVSgqbTzuqIAi9u7t_3WfP_3E_e7j7-Z83T5lBAaFnwr-WpbB5UGpT6lzAQdiPFu_d8zwvWPGoKEgcoAJbWQwsSZuPxv14VKXQVq0XF_zxJ5Cib5lstOJ0C3BoEqAUhh005MPQRMiQJTxiijkrzjGdV5yKtn4Qz-Va0ArguBX2kGOwkoaaY4F8U8nwGA-c0yCITZxBWJV5MW3kNOwgIsy8wMd-QIih418mQ5N8hZEhE4wwY4THXkJ7Ch2CaZHe_ny2x8dh1aTCXG63JtJ6Nn1e_NPkJ06NOQ9aFxuqYCFFo493ehYnHvGBl4bjIhKSFF-KuJC9nOilRLznFVjTRxcjihHOE_41I2OFWy5SwW62pV6ktz0eYWgsHo9ps0J4FX3T4GlVzWhTFrNKFq_NwOjLfZJ3Jg73Ivg-JsBLk1EZSBzbWPcuQ-VRTpkNIIQumSxnOOA46WmzfDS1b1_6XMq14kHsxT2rQ0965CWmKCE8IpiPse-bJg2RESiGwzziPdEqvqbFR1ciPoX8ua2rg6AD8e5tyR4IF5kQQv7FohQwgknYbVsiJkZmtHjlgj20um41wHFve1HStZCNLovmRqxLwbuq1AP5wSnj6pmbA_KhNvZqek4fgHah-QWUayG_ALqjYn0OpBe42rRGT3w1XN5bLT0i2Bkh7TzeR_kptGqbcsef-Gou9KNWgMw_t3errsYuy4o_fzSabwG5MSQazdxOK8r_bnlWH2Gaal-_qVLzbF_ze0iMr2hb6Yzximv-BSK56XieqjSXomm35qRPP2sNMnUNBSN1Pg9DRIoxes79VLCKvygqGteA_iHKoyBKn1_mUmj-rgfBGpgrEQovi8BJWCQk6YvwaLqh9PllyOCZb2nn_1zK6vPdsUKRTdpoRFuOCkYHrJbKyHyICJct77ywV6ZDRcWezfSRklqwhNEwGeaBade5svcdUfAB8X0Expdl-bQiduVopJoWnLEiyQc57xh9SrGTBT1leiJlL35sDcfepUJDaYH8OIdFgbItLYd1K1WKfpi8JPPhIdwF8PHBvJMl6_H2I9_ojNFIc0HzhObUQI0U6iLPA0SZgRop1KbfSeL-iTg3l5cnXsgdVx_7wLY7T614pitefQz03fe7mbI6G3WgiyCrcddpmMszHmmYjLQxM8GGR7M2ihChfW-7O0rfpfQkoOfdnXbg_yNEA3CpGo-ApoPETNyyS0t7WuHw4mmVR1FOrK_HMjePGKHoVzX-h8yby-r-Q-Z3ZaMvnSxbWlX_4oWW3ek8P9TsmpamxNsW8vsW6or3JTt5oWejY6zI5DFLVsUljove70H2n0n7S9b7f7RClxKhOaPwSG2jNOd53mm9LzR7rhvqTrx0KPaLlFWXtCe9aUyQaTvJSNtJEx4mAYWG2d47BgVdugthigpTiQi-DIHjHGaZ2c4aTZXOPqnbukBG6hOleRyxADpMZ7jodIbElZJ2puamhqvuQNlPs-CKlhVn9ooN-XtpbMO4YW9KVdtwONmZu2_DOSx750Xy71nk__ks8siTQmpLkR1ac7ht7HDQziD_xmOHnN0MMH15Sec_4MuPG7h8uLt7-PPnP2_h8ufdzTN8ebDLsz9u4dPN48PTi5uFPJ4b4MWNya2LQ7wDgOKwMkHFGaTa0fu-2niptzXAy_3katIN6Saxh1dePt2UZ4aqf4Nes_lVr_xtx3TpqSXkXQtqE-rXB5r_KxPNXxppHo3SL-26P_VHzRlf2ZnUYCIDm_IvnmlAZj2oA0CzkUob5lludNMoHAL2dwZ_xDDLqNaqzFvNs8y2ZTF_19nOHmyZIQBwbJs191i6jabGw6WoSsHPkXCWcShHDLf0lWeDJVMk-wK-mxNrsPJxsvLXycrbyQo9WclPVoqTFWZPw2gGFdetEvanlcgKYflaVpa6JWhpMAiiBSAz-zLxThhfOev8eX-fZijsnizLV4cXzLK8LStdioxu2boQmYuPlYV9t4_7643gzIYq8FL-bofMAOM5wHjM_vP9GW8d4WIIZpksXqtszXVmik6VmftFfDT_c47-mtO6kjmtHAPb2p-BMarVbbaqqM7epHrN1kq2dWbEAThGRkyCDyE2MhAfxAnAKWRUU-_YhfsNdGkDdy4GXgptclUle_cgIIsLtvGcQaA9w1MYQ1uYlgN8NIKPhvj4HD4ewcdf4NOqdeKbi5fRxTWT-ODHAxgaguEQncLgPcyQmoU8i0COEDAgc0DmzqWzbhG5xeiwaIR2a2F_rYPz-2sdwaBfg-AwbXe02psgNQng_D8z0p0K7B8J3PEkZ2TDZ-RwlphZYU9oB-eNER5w_BEV0JEK2KEEnyjHGPiM0k5YAvAMjVuMjCLjcWR_FJmMIwejyPE4cjiKPDipTnCjcXt545zjcewvzJ2c8a7DDgx2OCo4Og7vITb5wlvoXGh9ovtfmBwdx9mR7DgeF_440o7Qff8L9sfBdmy6L_yGjuPtCD8a2v64s6oVL2ijvV5qDyEsduy1I1IcEUgdmaPjyZXYPbvL1GLbZe6PrJEGA-D4bMNlyiXAc1NA3Iu4l-9egXuF7hW5l3Px_HwHZ4zswBL3Qh191C2jjhHqOKGOFQrGyHZ2mP_mOzmtz-H5-hydrc-fAeO5TTcrIQtnfxDMPBAszkO78Ap7mT1ARN9A3Cf1ABF_A3GfzgNE8g3E0D-D6H8DMfbOIAbfQEzOGSf8jlXROetEB8yuER9ev67YNWEJSegVv0aRH_jY9xNytbmOc0RQEhDKMKKckFUYo7jIUZ6zoAhRfFVeYw8Hno8JCr0YkWlAghWjnBHKwgIVMfA9vqVlNTWXz6lU6yv7AdA1IpHnxVcVzXnV2E_qMBb8DdpdgI07r9S1nfvk7boBvleVjW4-yehSV_zafVrlvhyS4jcuDletqq5__yOmTo3dNf6fAAAA__-nxMDU">