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

    <tr>
        <th>Summary</th>
        <td>
            Instantiating a CUDA kernel with a type predicated on __CUDA_ARCH__ causes cudaErrorSymbolNotFound
        </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>
    Example:

```
#include <iostream>
#include <typeinfo>
#include <cxxabi.h>

template<typename T>
__global__ void SomeKernel(T) {}

template<int>
struct SomeType { };

int main() {
#ifdef __CUDA_ARCH__
    using T = SomeType<0>;
#else
    using T = SomeType<1>;
#endif
    SomeKernel<T><<<1,1>>>({});

    if (cudaGetLastError() == cudaErrorSymbolNotFound) {
        std::cout << abi::__cxa_demangle(typeid(SomeKernel<T>).name(),0,0,0) << " not found!" << std::endl;
    }
}
```
</pre>
<img width="1" height="1" alt="" src="http://email.email.llvm.org/o/eJx8U92OmzwQfRpzY21kxiTABRdZsvm-qn9SN71GBg_ErbERNrvJ21cGstvVto0mifDMOZwzMxbOqc4gFmR7T7aHSEz-bMfia2eUMJ-j2spr8XAR_aCR8D1hc-zYGmxPgCvT6EkiJbxU1vkRRU_4w7ucvw6oTGv_lGsuF1GrzXnNsb3HftDC44ozokd6WrJV1WlbC11V9MkqSR9tjx9xNKgJZCcCOSXpPUkP74iU8QuD8-PU-Bl4ug4Y6mkA8PsFo4ynvVCGQHajWwS3EltaVeX3w77afyv_ryrC9pRSOjllOnqihB9eaAkvWXgfX9GoHf67PH5TbqRq1_rfLPJybgMvl4gJrLA5IFu9Q_7iJhColhLImkmK_9B_Es4_jKMdb_74IQgJ2fn48drXVn-x_mgnI187QNeP8zKsAt83dvJ00UFFrZbDqmouopLYC9NpJJDNc5cEsvcuIN-EyS46CJTs9ZvfiAkANdbTdhETh-c19SIEjdSL36Bvnf3ye1vVSBZc5jwXERZxynewzbMdi84FtimmrZDItsA5Z3XNWb5LY4h5zbOURaoABlsWszROIInzTZzWSdImLGuybZKlCUkY9kLpjdZP_caOXaScm7CIIc4TiLSoUbv5fgEYfKZzlgCE6zYWAXRXT50jCdPKefdK45XXWHwwzgvjlfBhaQQN20d_zn2kz8qfqaChxXQYUapGeJTUmrdbShsxOXR_G3E0jbo4ez-40E44Ejh2yp-netPYnsAxCFr_7obR_sDGEzjONhyB4-rzqYBfAQAA__-YXENd">