<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/68478>68478</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
[CLANG][GPU] Error using printf with -fno-builtin[-printf]
</td>
</tr>
<tr>
<th>Labels</th>
<td>
clang
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
Maetveis
</td>
</tr>
</table>
<pre>
# Description
When `-fno-builtin-printf` is in effect (such as when passing this option or simply `-fno-builtin`),
compiling a device function with printf results in an error.
# Impact
`/Oi-` commonly used on Windows when debugging (this is the default setting of Visual Studio) is translated by `clang-cl` to `-fno-builtin`.
There's no option to re-enable a builtin that was disabled by a previous flag, so the only option is to not use `-fno-builtin` or list all builtin functions except printf.
# Environment
```console
❯ cat /etc/os-release
NAME="Ubuntu"
VERSION="20.04.6 LTS (Focal Fossa)"
ID=ubuntu
ID_LIKE=debian
PRETTY_NAME="Ubuntu 20.04.6 LTS"
VERSION_ID="20.04"
HOME_URL="https://www.ubuntu.com/"
SUPPORT_URL="https://help.ubuntu.com/"
BUG_REPORT_URL="https://bugs.launchpad.net/ubuntu/"
PRIVACY_POLICY_URL="https://www.ubuntu.com/legal/terms-and-policies/privacy-policy"
VERSION_CODENAME=focal
UBUNTU_CODENAME=focal
```
```console
❯ clang --version
Ubuntu clang version 17.0.2 (++20231003073124+b2417f51dbbd-1~exp1~20231003073217.50)
Target: x86_64-pc-linux-gnu
Thread model: posix
InstalledDir: /usr/bin
```
```console
❯ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2022 NVIDIA Corporation
Built on Tue_Mar__8_18:18:20_PST_2022
Cuda compilation tools, release 11.6, V11.6.124
Build cuda_11.6.r11.6/compiler.31057947_0
```
```console
❯ hipconfig --version
5.4.22801-aaa1e3d8
```
# Reproducer
Assuming a file named main.cu with the contents:
```cuda
#ifdef __HIP__
#include <hip/hip_runtime.h>
#endif
__global__ void a() {
printf("my name: %s\n", "asd");
}
```
For cuda:
```console
> clang -x cu -fno-builtin-printf main.cu
ptxas fatal : Unresolved extern function 'printf'
clang-17: error: ptxas command failed with exit code 255 (use -v to see invocation)
Ubuntu clang version 17.0.2 (++20231003073124+b2417f51dbbd-1~exp1~20231003073217.50)
Target: x86_64-pc-linux-gnu
Thread model: posix
InstalledDir: /usr/bin
```
For HIP:
```console
❯ clang -x hip -fno-builtin-printf main.cu
main.cu:5:17: error: unsupported call to variadic function printf
__global__ void a() {
^
1 error generated when compiling for gfx906.
```
(This happens with both with `-mprintf-kind=` `buffered` and `hostcall`)
OpenMP offloading is probably also affected, but I didn't try it.
# Analysis
Here's what I think goes wrong:
Clang rewrites `printf` to `vprintf` for Cuda, and directly emits runtime calls for AMDGPU, but this only happens when builtins are enabled, when builtins are not enabled a normal function call is emitted to `printf`, but it won't be found because:
- The cuda device libraries don't provide it, as nvcc normally rewrites it to a special `vprintf`.
- AMDGPU doesn't even support C style variadic functions so trying to codegen the call to `printf` fails earlier (at compile time instead of link-time).
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzcV0tv27ry_zTMZiBBot-LLBw7PjX-TROkSQ-6EihxZPNfmhRIyo_N-ewXQ8lxHm3PuTi7WxSOzSHnPb-ZEd6rjUG8ZqMbNlpeiTZsrbu-Exj2qPxVaeXpmvEBLNFXTjVBWcOyJcvm3eefWzTAxllSG5uUrdJBmaRxyoSajTNQHpQBrGusAjA-9W21BeHhQM8aEm42ELbKg42swTrwatfo03umbJwxPmN80cmt7K5Rml4LkLhXFULdmiryOKiwhU4HcOhbHaIWwgA6Z136Wn8ybb1rRBXenJKw1b1KyIbK7nbW6BO0HiVYA38qI-2hN0Ji2W42pAjj02iJ8hC2CBJr0eoAHkMgsq3hm_Kt0PA1tFJZxmfxqhPGaxFQQhmNrrQwm6TSJDrYn7jhjf5PW3TI-MSDsWcfBgsOEzSi1AgC-pcQtiLAQXiQyhMpChTQONwr23qotdgwvgBvo_7R5J4j6WnB2EA--IlKFDatfACh9Yu8czw84LHCJvQh-eD-W7NXzpodmg8xiP8ra7zV2J_ecjZbsvkKKkEZtcJQMb6yPnGoUfj-2pf53S0bLBnnz2VrQss47wjfbh-_ru-_dDSepdkwHcPnp68UvZWthIaV9V7EVOufrJdssGw7Nv1B8Xn9f8RfYqlEXw8Pj7dPT9-L95LhlZD3WhSR9VmRF-qn-7vb4vnxc0fbhtB4NpgzvmJ8dTgc0k6XtLK7eNa_-vr88HD_-PSrh1vUzS9e3jz_UTze_u5x2W58qkVrqm0jZGowML7qXfKK0cPj-tt88b14uP-8Xnz_xyZo3AjN-Cqg2_lEGJk0VqtKoWd81Ti1F9WpOzp98ODifnnbu7ym8HXU55vnL0_PvyC-JNZ_lW1UlZAke3T-BQL7CHe0ngL5JM1STvnE-A3jNzzjgzzLBtlkkPMh4zclH-aTepTLspRJ_hcem_yvV5d4PklHEeu6Ahdug4EN5nCcjovxMGmqRCvTHpONac8Y4FBI2FmJmi421qtjn6rGB6E1yqVyRKKweUcRVeZfeMPsq-q9M-iMRHz5tl6uSdL0kSBu0UoBHVqjA-nUHl33YGGbk1ObbWwMFd3lWTZKeMb5mcnCusY6cWk6N4QthMFPLRZ3whXFtMinbDCPHzwrHr4-FcShF3ERLnpotNoTyPVoAXmejun3N_qSUoBe5EioWimKSHDdvdXZkHSQZ6PJbDgpsn_hxq1qKmtq9SGxRukw5Xya5YkQIseBnP5OCh_AIzbOyrY6-7b7nHvf7roeWSuNYMQOJeyEMmnVdm2SkL6yJqAJsULf6d1K8SJE1RJrKIpP64eiuJyaSrcSgQ0WW9UQ0KimcK0Jaofplg1uX26ikap-rV5RbLQthS4K2FslQcSamQGb3HQXAKBvGpHCd6doQpfII89GC0OAwBfAOBdexh8zNuifs8nyN25bWRcD_BOr30RrcHsu_iNULfxk0Dl7tHvQhKPwUIsgNACQrs_Gobd6jxLwGNBdWiMwPjkbOOkHm9j-8wk9jONKrOjIkwYRYSTUQlH3jvHDowpQWYnARyMqJGrQyZ7atUcEZfa26urnDCj_U6AF7yP6af3wdwF9j-lHqsO_j-v5x2A-Irx5G6DW-LZprKMxrqIhKFjYC6eEVNUl2n2o_3nyv_7HRn0p5Z1Y2KBBFwfHOIheJuKaiPVxlo3T38LG9InG1a1oGjS-S6fShm33jYa8Xadw8kMZSZ2c_D3Oyrau0aGkn5SObJxtrQ9kdj-kv5Jy36C5ewBb19oKSdopD42zpSj1CYT2FkTcDlBSHZdtgDVIJQ3jkwDBnUCFDxPj3Ah98sq_Pv90noMPNOauaakwP2Bj0cPBWbO5ZEXXGWLsHR6cCujJhsvS0k3d-8sBOZRaCWlIFkvlsAr6BLhTwUOPdjHwPl6e3y3_eHg-G9TtNzROv_iaAtYnmwfhELppPfrgI5EG7_4CCDDW7YS-ZFXMN-WjMpQNnfov2p-1UAEOtnNriVDb1kgosRKtx3e-SeCJugL1zn610qp0win0IHsWjbN7JRFUiE7x3UjQqaZPF8eqQPoI8A1WSui3jk3P8jp_gbToO_a4RwN9ScECfDhp_FhQPu4q7hSXSBthcIOm62l9Eb4JLAGnBxROK3QEdiKcRxOIEVTGB0IlW4NW5kdCh4zP3iQg9JhzJa8HcjaYiSu8zsez8XgymkyHV9vrKhdjkZe8RinEUI4nvJpm0_F0LJHP8mx2pa575Jzk2XA2HKa1mJVZjnmdibLMeMaGGe6E0qnW-11q3eZKed_i9Xg6nEyvtChR-7iwcx4xjPreaHnlrul-QsM6G2a0j_kLh6CCjlv-4vP8yx9stGSjG8rR0RJuI5y0cRfvkS9CwJslb3RzXutHy6vW6eu3Q_1GhW1bngd6vT__SRpn_x8r2heiDTTSRzP-EwAA__-HqP9S">