<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/131749>131749</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
[NVPTX] `sqrtf` always approximated with `llvm.nvvm.sqrt.approx.f`
</td>
</tr>
<tr>
<th>Labels</th>
<td>
new issue
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
Lai-YT
</td>
</tr>
</table>
<pre>
I've encountered an issue with the square-root (`sqrtf`) floating-point operation on NVPTX. It always compiles to the `llvm.nvvm.sqrt.approx.f` intrinsic, even when `-fno-approx-func` is specified. I have a hypothesis about the cause and would like to explore a potential fix. Any guidance would be appreciated! 😃
---
## What's the Problem?
Here's the relevant code:
```cuda
// cat main.cu
__device__ float f(float x) {
return sqrtf(x);
}
```
The function `sqrtf` always compiles to the `llvm.nvvm.sqrt.approx.f` intrinsic, regardless of `-fno-approx-func`:
```llvm
; clang -Wno-unknown-cuda-version -fno-approx-func main.cu -S -emit-llvm --cuda-device-only -o -
; ModuleID = 'main.cu'
source_filename = "main.cu"
target datalayout = "e-p6:32:32-i64:64-i128:128-v16:16-v32:32-n16:32:64"
target triple = "nvptx64-nvidia-cuda"
@.str = private unnamed_addr constant [11 x i8] c"__CUDA_FTZ\00", align 1
@.str.2 = private unnamed_addr constant [17 x i8] c"__CUDA_PREC_SQRT\00", align 1
; Function Attrs: convergent noinline nounwind optnone
define dso_local noundef float @_Z1ff(float noundef %x) #0 {
entry:
%__a.addr.i = alloca float, align 4
%x.addr = alloca float, align 4
store float %x, ptr %x.addr, align 4
%0 = load float, ptr %x.addr, align 4
store float %0, ptr %__a.addr.i, align 4
%1 = load float, ptr %__a.addr.i, align 4
%2 = call float @llvm.nvvm.sqrt.approx.f(float %1) #3
ret float %2
}
; Function Attrs: convergent nounwind
declare i32 @__nvvm_reflect(ptr) #1
; Function Attrs: nocallback nofree nosync nounwind willreturn memory(none)
declare float @llvm.nvvm.sqrt.rn.ftz.f(float) #2
; Function Attrs: nocallback nofree nosync nounwind willreturn memory(none)
declare float @llvm.nvvm.sqrt.approx.ftz.f(float) #2
; Function Attrs: nocallback nofree nosync nounwind willreturn memory(none)
declare float @llvm.nvvm.sqrt.rn.f(float) #2
; Function Attrs: nocallback nofree nosync nounwind willreturn memory(none)
declare float @llvm.nvvm.sqrt.approx.f(float) #2
attributes #0 = { convergent noinline nounwind optnone "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="sm_52" "target-features"="+ptx85,+sm_52" }
attributes #1 = { convergent nounwind "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="sm_52" "target-features"="+ptx85,+sm_52" "unsafe-fp-math"="false" "use-soft-float"="false" }
attributes #2 = { nocallback nofree nosync nounwind willreturn memory(none) }
attributes #3 = { nounwind }
!llvm.module.flags = !{!0, !1, !2, !3}
!llvm.ident = !{!4, !5}
!nvvmir.version = !{!6}
!0 = !{i32 2, !"SDK Version", [2 x i32] [i32 12, i32 6]}
!1 = !{i32 1, !"wchar_size", i32 4}
!2 = !{i32 4, !"nvvm-reflect-ftz", i32 0}
!3 = !{i32 7, !"frame-pointer", i32 2}
!4 = !{!"clang version 21.0.0git (git@github.com:Lai-YT/llvm-project.git e57cd100ca297cf81854e35cccbf703edddd4aad)"}
!5 = !{!"clang version 3.8.0 (tags/RELEASE_380/final)"}
!6 = !{i32 2, i32 0}
```
## What's the Possible Cause?
As defined in `__clang_cuda_math.h`, `sqrtf` expands to the `__nv_sqrtf` intrinsic:
https://github.com/llvm/llvm-project/blob/2f808dd0702ba3c364eb4373714a1fb09078909d/clang/lib/Headers/__clang_cuda_math.h#L318
I noticed that `__nv_sqrtf` is defined by NVIDIA in _libdevice.bc_, so I grabbed the libdevice on my machine (which is `libdevice.10.bc`) and disassembled the bitcode with _llvm-dis_:
```llvm
@.str = private unnamed_addr constant [11 x i8] c"__CUDA_FTZ\00", align 1
@.str.2 = private unnamed_addr constant [17 x i8] c"__CUDA_PREC_SQRT\00", align 1
...
define float @__nv_sqrtf(float %x) #0 {
%1 = call i32 @__nvvm_reflect(ptr @.str) #6
%2 = icmp ne i32 %1, 0
br i1 %2, label %3, label %10
3: ; preds = %0
%4 = call i32 @__nvvm_reflect(ptr @.str.2) #6
%5 = icmp ne i32 %4, 0
br i1 %5, label %6, label %8
6: ; preds = %3
%7 = call float @llvm.nvvm.sqrt.rn.ftz.f(float %x) #6
br label %__nvvm_sqrt_f.exit
8: ; preds = %3
%9 = call float @llvm.nvvm.sqrt.approx.ftz.f(float %x) #6
br label %__nvvm_sqrt_f.exit
10: ; preds = %0
%11 = call i32 @__nvvm_reflect(ptr @.str.2) #6
%12 = icmp ne i32 %11, 0
br i1 %12, label %13, label %15
13: ; preds = %10
%14 = call float @llvm.nvvm.sqrt.rn.f(float %x) #6
br label %__nvvm_sqrt_f.exit
15: ; preds = %10
%16 = call float @llvm.nvvm.sqrt.approx.f(float %x) #6
br label %__nvvm_sqrt_f.exit
__nvvm_sqrt_f.exit: ; preds = %6, %8, %13, %15
%.0 = phi float [ %7, %6 ], [ %9, %8 ], [ %14, %13 ], [ %16, %15 ]
ret float %.0
}
```
The selection of intrinsic respects the value of `__nvvm_reflect("__CUDA_FTZ")` and `__nvvm_reflect("__CUDA_PREC_SQRT")`. The [_NVVMReflectPass_](llvm/lib/Target/NVPTX/NVVMReflect.cpp) is the pass that picks them up and replaces the `__nvvm_reflect(...)` expressions with appropriate integer values.
However, it only recognizes `"__CUDA_FTZ"` and `"__CUDA_ARCH"`, leaving others with the default value `0`. This is why `__nvvm_reflect("__CUDA_PREC_SQRT")` always evaluates to `0`, leading to the selection of `llvm.nvvm.sqrt.approx.f`:
https://github.com/llvm/llvm-project/blob/2f808dd0702ba3c364eb4373714a1fb09078909d/llvm/lib/Target/NVPTX/NVVMReflect.cpp#L168-L178
I guess that the old libdevice doesn't rely on `"__CUDA_PREC_SQRT"`, thus _NVVMReflectPass_ doesn't try to handle it. (The last commit on this is 7 years ago. 😏)
## Any Possible Solution?
_Clang_ provides several options that affect the behavior of floating-point operations, and three of them are said to take effect in CUDA code. (Stated in [Compiling CUDA with clang, Flags that control numerical code](https://llvm.org/docs/CompileCudaWithLLVM.html#id5). Some of the flags are renamed after then.):
- `ffp-contract`
- `fgpu-flush-denormals-to-zero`
- `fgpu-approx-transcendentals`
These flags don't seem to handle `sqrt`. Maybe we can add a flag, e.g., `-fgpu-approx-sqrt`, or respect the general `-fapprox-func` flag? And set a metadata just like `nvvm-reflect-ftz` for _NVVMReflectPass_ to pick up. Or, look at the `"approx-func-fp-math"="true"` metadata, which is already set when `-fapprox-func` is given.
To add metadata, the following part is probably where we should place the logic:
https://github.com/llvm/llvm-project/blob/1fbfef9b8ac684483619811d9e90cf8e1f84c6ee/clang/lib/CodeGen/CodeGenModule.cpp#L1289-L1296
---
To sum up the issue, I believe the non-approximated square-root (`llvm.nvvm.sqrt.rn.f`) should be used by default, or unless there should be an option for the user to control with. Just like _NVCC_ provides the `-prec-sqrt` flag, _Clang_ could provide this as well.
If this is a missing puzzle, I would like to work on this; if it's already supported or not desired to be touched, please let me know!
</pre>
<img width="1" height="1" alt="" src="http://email.email.llvm.org/o/eJzUWltz47hy_jXwSxdZJChK1IMfZHmddTK72exMZpPzogKJpoRjCOABQMmaX58CQOpiy56Zs1PJZso1urC70ZevG92AmLVirRBvSXlHyvsb1ruNNrcfmEj--9NNrfnh9pHQ2Q4BVaN75dAgB6ZAWNsj7IXbgNsg2H_0zGBitHZAaEWmmf2HcS2ZZoTOoZWaOaHWSaeFcqA7NMwJrUAr-PXzb5_-K4VHB0zu2cFCo7edkGjB6SCaTDMpd9tU7Xbb1EtNWdcZ_Zx66SCUM0JZ0RC6BNyhgv0GlWdKWqWTSJq0vWoCtQXbYSNagTyFR9iwHQKDzaHTboNWWGC17l1Yt2G9RWCKw173koMUT-h1wudOauPZOu1QOcEktOI5hYU6wLoXnKkGB54awWuAjWAOOaE5kIeMzB_IvCJVASRbkGyRJEl8Q2hBaAF_bJgjdGaDFr8ZXUvckuIh0vyMBo8PDUrcMeWg0RxJsRjETLP41_ScBbEPhD5AwxxsmVBp05NssVpx3IkGV6sYHWgJreK7Zx8yMrsj2QLAoOuNghhNWvlnpPCPyOz-fK249KcNgvd1CO4ZCH5AbA2umeESrQXdvhHf1y7w0v3n4g4aydQakj-UTnr1pPReJd5ByQ6N9eq-lDc6C5KPkOBWuMQLgyRyRe8lWskDJBqSYZFfNO8lPt4DKe6B0NnocDoj2cLq3jS4aoVExbY40NAjDSXZwjGzRgecOSbZwWNxoMKkm5JiUdDwXyKmE1IsppNE5LQixSKnVbLLPUU-TXYjlcqPPNPJxQLOiE4eVVC7zj1PJ4naCS5YsDBS-79JllpnAmlnxI45hF55A_iKcW6g0co6j0JS3uU5PIOoSHkPDaF0tVr-5_1i9fDpb6RcZpmXSZfApFgryE-yU_qN0mfXpP_2-0_L1cf_-P3TG2vEyDyMsFw4ZywpFl7yDs0alQOlhZJCISjdq71QHHTnlFZIsgXH1j_hVq-kbpgMNBzbIW_IJFv9LW9P6TM-JrSMmUSLbEgnVM4cIkjBP1-tWOqNTEWwn0m_QJR7smEyUj8H2q9SWuer06Bc0GEJnY_fKOGK6CxIlZrxk8x3eV4skp0xnIy6slD-9kLv8kWANEzKk9vfqhxjIPxyQwCKsZadVKanGvZNAIm4CHhoJDMIoqAh-CuvwspgK7FxhFadM8Oq76JPeSzJmjVPoHRr0GPPHlRzguBeSDmU3y1utTkQWgVM0vmZGm-5w6i0dV9O7hh0ov-XOo0h-qvp5X31V9LnJZQvdWLOGVH3Du1QW3wNn919Uz3zxb41bIuxFUPjy2VxTyhlUhJKw2agE2dY1_mGbcvc5kjiTI8DjXWseUo6ox02Tpuk7tsWTWLFFzySVwNt3HGSpuuPj-x2VdLLxy0y1xu0RxpC7zr3XJWELgm9O3GEnL30QX7dB4PphFIuLKslJo4JmfjYnZZpmbSjVV_zjO8-ktDOWUzabsv4VTFKJ0K1Nmm7S_-9oFFMfZXmrxwISntlWYvv2tBbTKxuXTIA-RXJlXDSYzj_VKJdlV2cyR7xcdwFaB4ScRuauLSVbG2HDin3GzjNwz5HaJ4Pr3R4LQYZgwDBPQAvOCcDZXmk9OkuTDq2nxfU03OdsrNnftMZFyWUfrz_N_gcBQx9DynvqG-SCuq7JFIGjqinfzcl5f1Rg_yF4PwkeN9smFmNKBqYJ0dO-oJzcuL0ZiXDbpi07ssZf3bkL17wz078r3JwYKZH5smlrwilsbUfPUnzNEuztQiT6Fo4MvEfNn2dNnpLikUcbgl98LHyqfN3bFzq6bGcNTzPsobR-axpq7wqJ1iUTdPU7SwrkHPOJ8xn_dxrNipUfkWhIq3SzCvj2NoS-vD7Tx9-Wnz8aVVUGaEPrVBMvpA4vRbycxdejl1XB0dtraglwtLPscf5cWEhNrQcRBjRVqug7Mr3_CufxukmTO3Li_kNnzum-Pnc5vue1fH5aVQbR7CNc53fLOP0eRaA6PcX7if0oZa6JvSBtlVWcZ7NMlqzoimmE6wnxayY5ROWt3U2z2bVPJtzQh-C4l6Q8Iw_I-NovH-vmUSLD0VeRdUeQWknGuTgNn7_fWXMyUf1AX79_Hj_uPDeWklRx7kvrZuVd5HV8Ahrw-o6CEM4UoBWsD3AljUbEbbdar8RzcbL9mPpUVCepXUzHJMwxcHvVNbitpaDxFo4P93Ho5ZV8BkXdvX2qPv_f1xL0zS-GUav06B1CtNZk_9yyjobNMLA8E6jDoNFg4Tp5bghmm0Hauj0wzCxhCyQ1AZEHqcIugTJapT-U3HxKR-ys_At43f-831nZ5CP-0-ZjapNvsuwlL40rbxm2uSKaeWFMdOLT0MiTX-EZcWo2ezrM97LoeY8_NNR_aOWg2M856pN8Vm4qHb1Q9Wef_to-udVz7Pv1_1tLOXflyWvwZRfT5RrmZJfpkr-IlfKwb5_Ille2ZefDJx8G6b-ZFDKH6v09J877Ph-xa98_1VDXik-jb1bWQ2vMbJjTINFaWxhu40YDSrvQsYPlFPwbWnsXkNKjSJffD920aVvIC-fTI-LhievjnvS7L0za4se7eFGoj31MmDQdti42FDtmOxxOHt-lSWXm6ff0-bh2NuPF--Sn3bDkSkFrxEp71a_fv78y--R5zdm7SqYXI3dU2h6PoWZjdCHcI0SXo88adN1HhIi6t8xa2PD04nmKXy3hb4LOhrsJGvQnjd3Fwr7LTmahM-dQev7Wht7kgDFzgjfCvimfY0m-sr6XfxnvccdhvND4SAclxts9FqJLxhaoVeeO7nt9Gjx-_Ln-CxUDWQ7odag3QaNPV1CcWxZL90QKTLNssGdwnon7DeH7w7GeHGBXiZz8eZikDyowr0qQ198gaN3Lzf-N_vk70RM8SGfVsmHfHZsl9c9jtjxVupwGzY2ulyjVYTOHBiUB4gXP2_4NDrNbXoLr9B9JsiZg_fohikuEYRLff_ss0Iy66DR223AErghsjM4IDMW2Fqn5zdsD_G47TQhLdThNBh91LL3sTrORqtlGBugM3onOFqwHrpMgu5cwHvwAGtbbKIjatywndDGR_utO04b2lvlm3mDoX6ExGMGwTLBA3LYEwJGsUKBd1u40Qtmf3TMDcNaebcM92cecIEoIH-YgpbwEM4rgo6NVs5oCarfohENk_GGMNSPS8AFhGrjxyiuGz89xTVw2XP2h3CbDx8-_5Ju3FYSWgheEjpP4aPejpZAPCXx5phwq8WBtQ6Nf6ZCzRiBnnhctG2XBOVY42IFjl-vuz5pZW83CUelzZZJmzidfEGjX5INV3TOMGUbVByVY9Kel3M7asV1xJNF3J4BahhsQ3H4hR1qhD1CwxQwzoEF3nCbnK7TYQ5OzlceeP0jbcY9IvhijSrgJXC8uHoOUosHWCgOFh0w2KJjnDkGf--ti9fLZJq9Oj_xvNpcyRenQyWHvkvh30OFlVo_wZCkMQnPlHh1TDecIvoFRlW8kOOcyqRBxg9B29Ol-usb9bXYoUqHy-xPOnjxXGBAiZZS7z1wO2ac5-qMrlktD160CRGwm3BnHnaiOEzr9Q84UMjbusV2XlesmVaTSVVM83mV53yO86xpK8zbatJMEV8dKCw1x39BdXoXL3ePRZJW8-RDTufT4zX-yQe2D3urtyL8UsL74RFqlAJ30Til1QAosQ0Z_vonFNca1XhSMLiqRuhtPKUY9r4BlL2SsWJ7156ImRpKWUCU16K3PlX1sWD4ipLCvx4Bufr183J5VhEHZIVD8DERjgkz1s8mxjHyxCLNLOxRymGuf2yPpZvBVlgbgNF_-SIHR13-4mKvzdNY7n37KVoQ8aDrCNG-67TxXtQGlHbA0QqDobrWXkbfbJCHy0aJzCJIdLBFeFJ6T2h-w28LPi_m7AZv89mE0qqYzejN5hbLeTVjGW_mE5rX9XyWTwvOZ9Mpy4q8nGY34pZmtMyKvMpmE1rSdJ7nVc5znNWsKJs8J5MMt0zIdCy0NwEPt3mRzybzm9Ci2_DrG0oV7ke0UFLe35jbgOi6X1s_AQjr7EmME06Gn-3Efby8v_Z7iwuAhd3i3abkpjfy9rvzLKjsd47Bpt0t_Z8AAAD__9Be-EA">