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

    <tr>
        <th>Summary</th>
        <td>
            [CUDA][HIP] comparison of device functions not allowed in host function
        </td>
    </tr>

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

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

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

<pre>
    HIP/CUDA has separate compilations for host and device. Instructions of host functions are generated by host compilation, during which the compiler has no access to device function pointers. The device functions seen by host code is just a placeholder address, not the real device function address. If the placeholder address is stored to a variable, then passed to kernel and called there, it won’t work. To avoid misuse, clang forbids using of device functions in host function.

However, if device function is used for comparison with each other, it should be fine, e.g.

```
__device__ float dfn1(float) { return 1;}
__device__ float dfn2(float) { return 2;}

template<float (*OP)(float)> 
__global__  void some_kernel(float *x, float y) {
  *x = OP(y) + 10 *y;
}

template<float (*OP)(float)>
void run_kernel(float* x) {
   constexpr float param = (OP == &dfn1) ? 1 : 0;
   some_kernel<OP><<<1,1>>>(x, param);
}

void run(float* x) {
    run_kernel<dfn1>(x);
}

```
However, currently clang diagnose the above code (https://godbolt.org/z/d4aW1oor4 ) whereas nvcc allows it (https://godbolt.org/z/YWYKeTr67).

Basically, nvcc only diagnose call of device functions in host function and allows other uses, while clang diagnose any ODR-use of device functions in host functions.

I think we may want to be consistent with nvcc regarding use of device functions in host functions.

@Artem-B 
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJycVUtv2zgQ_jXjyyCGRPkhH3Two0aDPSRYdFH0ZFDiSGJLkQZJ2fH--gUpuXWcZhEUIGST8_i-meEMuXOy0UQFzDfA2OXF8U7JHhiD-W7Ce98aW1xPJ6URl-Lz4zOw_faf3Rpb7tDRkVvuCSvTHaXiXhrtsDYWW-M8ci1Q0ElWNMVH7bztq0HD1INC3evxhFvChjQFbwLLyyC_cQtsi6K3Ujd4bmXVom-vsGQjGW2QVxU5h96MsD8B8Gik9mTdFL-0dC8NgZC-QRWE0uH3PsSAR8Urao0SZJELYcm5QEYbHzlY4uoN3Kg3xcc6Kv3GR0Bw3lgSgS_HE7eSl4qCb9-SxiN3bhD-IKtJxXRWXKlw2JKNmtLj2Wj4xCBPYLUKO_tjil8M8pORAjvpehc1K8V1E2pTSuGwdyGTpn6bCqlf12YKyQ6S9fD9bM50IhuR39iGiPpAOVyAUBpupTMaz9K3SLxq0QTeI23Xml4JLAlrqSNDmjavwGCRjCtuD4cB73DAWhnuUdQ6BZbHDbAVwnKDlnxvNaaQbWC5e9-QvWPIbg2Hr6fuqLgnyLaDObAc2PrpGdjqxgtkn_CK1yhTcnU4IMYiONPRYSji1QCBrV9C0MPuMtIYHGCUImQ7DCj5IGUbTJMguQSOA8E_ZTroR3K213fcgK3x5Y4QVkY7Ty9HOzIOjd9FisDyp-fwb9gshrKsELI9pgjZGpOfhBFfJSPbPj0HNtl2WCmwbRoP4mJ5TFGEirx_H_Y1jP_lfxtnto0crwjvOr67gDeXv-qtJe3VZWwrIXmjjaPY7Lw0JxqmCLC89f7oIFsD2wPbN0aURvmpsQ2w_b_A9mLGv6bG2BkGyufQ2GGYnaoKuVLm7EKzfMDPt6_f_qIvdrEEtnrVRhvuZBgblzi1gl-j1eUX5SD70CSIA2jkFDs5dHuchedWKrpPBdcXfNr9_dA7-pB794r1I_pW6h94Juz4Bc9c-zAKS4o3UTpP2g-TJYZkqeFWhJn2Z3AwS9bWU_ewwYkoMrHKVnxCRbpks1maZ_PFpC14nZf1jPF8kfJlJjJBdZ7SvFzmyaKk5WwiC5awWZKzLM2S-XwxzZJSVHmaiXq1TOpKwCyhjks1VerUhdpNpHM9FWkyz9l8onhJyo3PsaYzRun4HtsiGD2UfeNglijpvPvlxkuv4jsenmaY72C-CU_1fHc7h3-XlfCKxZKSeJOhSW9VcXftpG_7clqZDtg-oI8_D0drvlPlge0jZwdsPwZ1Kth_AQAA__-uXrEW">