Added detection support for pre-generated LibDevice PTX.

m4rs-mt · Jan 14, 2024 · b049269 · b049269
1 parent 8afc5d7
commit b049269
Show file tree

Hide file tree

Showing 10 changed files with 235 additions and 11 deletions.
diff --git a/.gitignore b/.gitignore
@@ -286,6 +286,7 @@ Src/ILGPU/Static/DllImports.cs
 Src/ILGPU/StrideTypes.cs
 Src/ILGPU/Util/DataBlocks.cs
 Src/ILGPU/Util/PrimitiveDataBlocks.cs
+Tools/CudaGenerateLibDeviceTool/CudaDriverVersionUtils.cs
 
 # Ignore specific template outputs (Algorithms)
 Src/ILGPU.Algorithms/AlgorithmContextMappings.cs

diff --git a/Src/ILGPU/Backends/PTX/PTXBackend.cs b/Src/ILGPU/Backends/PTX/PTXBackend.cs
@@ -14,13 +14,15 @@
 using ILGPU.IR;
 using ILGPU.IR.Analyses;
 using ILGPU.IR.Transformations;
+using ILGPU.Resources;
 using ILGPU.Runtime;
 using ILGPU.Runtime.Cuda;
 using ILGPU.Util;
 using System;
 using System.Collections;
 using System.Collections.Generic;
 using System.Diagnostics.CodeAnalysis;
+using System.Linq;
 
 namespace ILGPU.Backends.PTX
 {
@@ -405,8 +407,10 @@ private void GenerateLibDeviceCode(
                         StringComparison.Ordinal);
                 builder.KernelBuilder.Append(compiledString);
             }
-            else
+            else if (Architecture >= PTXLibDevicePtx.MinArchtecture &&
+                InstructionSet >= PTXLibDevicePtx.MinInstructionSet)
             {
+                // If supported, use the pre-generated LibDevice PTX code.
                 var ptxModules = InlineList<string>.Create(backendContext.Count);
                 PTXLibDevicePtx.GetPtx(
                     enumerator.AsEnumerable(),
@@ -416,6 +420,28 @@ private void GenerateLibDeviceCode(
                 builder.AddModule(ptxModules.AsReadOnlySpan());
                 builder.KernelBuilder.AppendLine(ptxDeclarations);
             }
+            else if (enumerator.AsEnumerable().FirstOrDefault() != null)
+            {
+                // Handle any issues if a LibDevice function is used.
+                if (Architecture >= PTXLibDevicePtx.MinArchtecture)
+                {
+                    // The architecture is supported, but is using an older instruction
+                    // set. Can be solved by a driver update.
+                    var minDriverVersion = CudaDriverVersionUtils.GetMinimumDriverVersion(
+                        InstructionSet);
+                    throw new NotSupportedException(string.Format(
+                        RuntimeErrorMessages.NotSupportedLibDevicePreGeneratedNewer,
+                        minDriverVersion.Major,
+                        minDriverVersion.Minor));
+                }
+                else
+                {
+                    // The architecture is too older for the pre-generated LibDevice PTX.
+                    // Inform the user to manually initialize LibDevice.
+                    throw new NotSupportedException(
+                        RuntimeErrorMessages.NotSupportedLibDeviceNotInitialized);
+                }
+            }
         }
 
         #endregion

diff --git a/Src/ILGPU/Backends/PTX/PTXLibDevicePtx.tt b/Src/ILGPU/Backends/PTX/PTXLibDevicePtx.tt
@@ -20,9 +20,10 @@
 <#@ output extension=".cs" #>
 <#
 string rootPath = Host.ResolvePath(".");
-var functions = LibDevicePtxFunctions.Load(rootPath, "../../Static/CudaLibDevicePtx.xml");
+var lib = LibDevicePtxFunctions.Load(rootPath, "../../Static/CudaLibDevicePtx.xml");
 #>
 
+using ILGPU.Runtime.Cuda;
 using ILGPU.Util;
 using System.Collections.Generic;
 using System.Text;
@@ -36,10 +37,22 @@ namespace ILGPU.Backends.PTX
     /// </summary>
     internal static class PTXLibDevicePtx
     {
+        /// <summary>
+        /// Minimum architecture required to use the pre-generated PTX.
+        /// </summary>
+        public static CudaArchitecture MinArchtecture { get; } =
+            new CudaArchitecture(<#= lib.MinArchitectureMajor #>, <#= lib.MinArchitectureMinor #>);
+
+        /// <summary>
+        /// Minimum instruction set required to use the pre-generated PTX.
+        /// </summary>
+        public static CudaInstructionSet MinInstructionSet { get; } =
+            new CudaInstructionSet(<#= lib.MinInstructionSetMajor #>, <#= lib.MinInstructionSetMinor #>);
+
         #region Generated PTX
 
 <#
-    foreach (var func in functions)
+    foreach (var func in lib.Functions)
     {
         PushIndent(2);
         WriteLibDeviceFunctionPtx(func);
@@ -54,7 +67,7 @@ namespace ILGPU.Backends.PTX
             {
 <#
     PushIndent(4);
-    foreach (var func in functions)
+    foreach (var func in lib.Functions)
         WriteLine($"{{ \"{func.Name}\", {func.Name} }},");
     PopIndent();
 #>
@@ -65,7 +78,7 @@ namespace ILGPU.Backends.PTX
             {
 <#
     PushIndent(4);
-    foreach (var func in functions)
+    foreach (var func in lib.Functions)
         WriteLine($"{{ \"{func.Name}\", {func.Name}{DeclarationSuffix} }},");
     PopIndent();
 #>
@@ -135,8 +148,20 @@ void WriteLibDeviceFunctionPtx(LibDevicePtxFunction func)
 [XmlRoot("LibDevicePtx")]
 public class LibDevicePtxFunctions
 {
-    internal static LibDevicePtxFunction[] Load(string rootPath, string fileName) =>
-        XmlHelper.Load<LibDevicePtxFunctions>(rootPath, fileName).Functions;
+    internal static LibDevicePtxFunctions Load(string rootPath, string fileName) =>
+        XmlHelper.Load<LibDevicePtxFunctions>(rootPath, fileName);
+
+    [XmlAttribute]
+    public int MinArchitectureMajor { get; set; }
+
+    [XmlAttribute]
+    public int MinArchitectureMinor { get; set; }
+
+    [XmlAttribute]
+    public int MinInstructionSetMajor { get; set; }
+
+    [XmlAttribute]
+    public int MinInstructionSetMinor { get; set; }
 
     [XmlElement("Function")]
     public LibDevicePtxFunction[] Functions { get; set; }

diff --git a/Src/ILGPU/Resources/RuntimeErrorMessages.Designer.cs b/Src/ILGPU/Resources/RuntimeErrorMessages.Designer.cs
diff --git a/Src/ILGPU/Resources/RuntimeErrorMessages.resx b/Src/ILGPU/Resources/RuntimeErrorMessages.resx
@@ -241,9 +241,15 @@
     <value>Unknown parent accelerator</value>
   </data>
   <data name="VelocityPlatform64" xml:space="preserve">
-      <value>Velocity accelerator requires 64-bit application ({0} not supported). Ensure Prefer32Bit is set to 'false'</value>
+    <value>Velocity accelerator requires 64-bit application ({0} not supported). Ensure Prefer32Bit is set to 'false'</value>
   </data>
   <data name="VelocityLittleEndian" xml:space="preserve">
-      <value>The Velocity accelerator supports little-endian machines only</value>
+    <value>The Velocity accelerator supports little-endian machines only</value>
+  </data>
+  <data name="NotSupportedLibDeviceNotInitialized" xml:space="preserve">
+    <value>Cannot find LibDevice implementation. Ensure that LibDevice is enabled from the ContextBuilder.</value>
+  </data>
+  <data name="NotSupportedLibDevicePreGeneratedNewer" xml:space="preserve">
+    <value>Cannot find LibDevice implementation. Upgrade to Cuda driver &gt;= v{0}.{1} or override LibDevice from the ContextBuilder.</value>
   </data>
 </root>
diff --git a/Src/ILGPU/Static/CudaLibDevicePtx.xml b/Src/ILGPU/Static/CudaLibDevicePtx.xml
@@ -1,5 +1,5 @@
 <?xml version="1.0" encoding="utf-8"?>
-<LibDevicePtx>
+<LibDevicePtx MinArchitectureMajor="6" MinArchitectureMinor="0" MinInstructionSetMajor="8" MinInstructionSetMinor="0">
   <Function Name="__nv_abs" PtxModule="//&#xD;&#xA;// Generated by NVIDIA NVVM Compiler&#xD;&#xA;//&#xD;&#xA;// Compiler Build ID: CL-31968024&#xD;&#xA;// Cuda compilation tools, release 12.0, V12.0.76&#xD;&#xA;// Based on NVVM 7.0.1&#xD;&#xA;//&#xD;&#xA;&#xD;&#xA;.version 8.0&#xD;&#xA;.target sm_60&#xD;&#xA;.address_size 64&#xD;&#xA;&#xD;&#xA;&#x9;// .globl&#x9;__ilgpu__nv_abs&#xD;&#xA;&#xD;&#xA;.visible .func  (.param .b32 func_retval0) __ilgpu__nv_abs(&#xD;&#xA;&#x9;.param .b32 __ilgpu__nv_abs_param_0&#xD;&#xA;)&#xD;&#xA;{&#xD;&#xA;&#x9;.reg .b32 &#x9;%r&lt;3&gt;;&#xD;&#xA;&#xD;&#xA;&#xD;&#xA;&#x9;ld.param.u32 &#x9;%r1, [__ilgpu__nv_abs_param_0];&#xD;&#xA;&#x9;abs.s32 &#x9;%r2, %r1;&#xD;&#xA;&#x9;st.param.b32 &#x9;[func_retval0+0], %r2;&#xD;&#xA;&#x9;ret;&#xD;&#xA;&#xD;&#xA;}" PtxDeclaration=".extern .func  (.param .b32 func_retval0) __ilgpu__nv_abs(&#xD;&#xA;&#x9;.param .b32 __ilgpu__nv_abs_param_0&#xD;&#xA;)&#xD;&#xA;;" />
   <Function Name="__nv_acos" PtxModule="//&#xD;&#xA;// Generated by NVIDIA NVVM Compiler&#xD;&#xA;//&#xD;&#xA;// Compiler Build ID: CL-31968024&#xD;&#xA;// Cuda compilation tools, release 12.0, V12.0.76&#xD;&#xA;// Based on NVVM 7.0.1&#xD;&#xA;//&#xD;&#xA;&#xD;&#xA;.version 8.0&#xD;&#xA;.target sm_60&#xD;&#xA;.address_size 64&#xD;&#xA;&#xD;&#xA;&#x9;// .globl&#x9;__ilgpu__nv_acos&#xD;&#xA;&#xD;&#xA;.visible .func  (.param .b64 func_retval0) __ilgpu__nv_acos(&#xD;&#xA;&#x9;.param .b64 __ilgpu__nv_acos_param_0&#xD;&#xA;)&#xD;&#xA;{&#xD;&#xA;&#x9;.reg .pred &#x9;%p&lt;6&gt;;&#xD;&#xA;&#x9;.reg .b32 &#x9;%r&lt;12&gt;;&#xD;&#xA;&#x9;.reg .f64 &#x9;%fd&lt;98&gt;;&#xD;&#xA;&#xD;&#xA;&#xD;&#xA;&#x9;ld.param.f64 &#x9;%fd13, [__ilgpu__nv_acos_param_0];&#xD;&#xA;&#x9;{&#xD;&#xA;&#x9;.reg .b32 %temp; &#xD;&#xA;&#x9;mov.b64 &#x9;{%temp, %r1}, %fd13;&#xD;&#xA;&#x9;}&#xD;&#xA;&#x9;abs.f64 &#x9;%fd1, %fd13;&#xD;&#xA;&#x9;{&#xD;&#xA;&#x9;.reg .b32 %temp; &#xD;&#xA;&#x9;mov.b64 &#x9;{%temp, %r3}, %fd1;&#xD;&#xA;&#x9;}&#xD;&#xA;&#x9;setp.lt.s32 &#x9;%p1, %r3, 1071801958;&#xD;&#xA;&#x9;@%p1 bra &#x9;$L__BB0_8;&#xD;&#xA;&#x9;bra.uni &#x9;$L__BB0_1;&#xD;&#xA;&#xD;&#xA;$L__BB0_8:&#xD;&#xA;&#x9;mul.f64 &#x9;%fd61, %fd1, %fd1;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd62, 0dBFB3823B180754AF;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd63, 0d3FB0066BDC1895E9;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd64, %fd63, %fd61, %fd62;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd65, 0d3FB11E52CC2F79AE;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd66, %fd64, %fd61, %fd65;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd67, 0dBF924EAF3526861B;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd68, %fd66, %fd61, %fd67;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd69, 0d3F91DF02A31E6CB7;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd70, %fd68, %fd61, %fd69;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd71, 0d3F847D18B0EEC6CC;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd72, %fd70, %fd61, %fd71;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd73, 0d3F8D0AF961BA53B0;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd74, %fd72, %fd61, %fd73;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd75, 0d3F91BF7734CF1C48;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd76, %fd74, %fd61, %fd75;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd77, 0d3F96E91483144EF7;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd78, %fd76, %fd61, %fd77;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd79, 0d3F9F1C6E0A4F9F81;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd80, %fd78, %fd61, %fd79;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd81, 0d3FA6DB6DC27FA92B;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd82, %fd80, %fd61, %fd81;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd83, 0d3FB333333320F91B;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd84, %fd82, %fd61, %fd83;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd85, 0d3FC5555555555F4D;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd86, %fd84, %fd61, %fd85;&#xD;&#xA;&#x9;mul.f64 &#x9;%fd87, %fd61, %fd86;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd9, %fd87, %fd1, %fd1;&#xD;&#xA;&#x9;setp.lt.s32 &#x9;%p5, %r1, 0;&#xD;&#xA;&#x9;@%p5 bra &#x9;$L__BB0_10;&#xD;&#xA;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd88, 0dBC91A62633145C07;&#xD;&#xA;&#x9;add.rn.f64 &#x9;%fd89, %fd9, %fd88;&#xD;&#xA;&#x9;neg.f64 &#x9;%fd90, %fd89;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd91, 0d3FF921FB54442D18;&#xD;&#xA;&#x9;add.rn.f64 &#x9;%fd97, %fd91, %fd90;&#xD;&#xA;&#x9;bra.uni &#x9;$L__BB0_11;&#xD;&#xA;&#xD;&#xA;$L__BB0_1:&#xD;&#xA;&#x9;mov.f64 &#x9;%fd14, 0d3FF0000000000000;&#xD;&#xA;&#x9;sub.f64 &#x9;%fd2, %fd14, %fd1;&#xD;&#xA;&#x9;{&#xD;&#xA;&#x9;.reg .b32 %temp; &#xD;&#xA;&#x9;mov.b64 &#x9;{%temp, %r2}, %fd2;&#xD;&#xA;&#x9;}&#xD;&#xA;&#x9;setp.lt.s32 &#x9;%p2, %r2, 1;&#xD;&#xA;&#x9;@%p2 bra &#x9;$L__BB0_3;&#xD;&#xA;&#xD;&#xA;&#x9;add.s32 &#x9;%r4, %r2, -1048576;&#xD;&#xA;&#x9;{&#xD;&#xA;&#x9;.reg .b32 %temp; &#xD;&#xA;&#x9;mov.b64 &#x9;{%r5, %temp}, %fd2;&#xD;&#xA;&#x9;}&#xD;&#xA;&#x9;mov.b64 &#x9;%fd16, {%r5, %r4};&#xD;&#xA;&#x9;// begin inline asm&#xD;&#xA;&#x9;rsqrt.approx.ftz.f64 %fd15, %fd16;&#xD;&#xA;&#x9;// end inline asm&#xD;&#xA;&#x9;{&#xD;&#xA;&#x9;.reg .b32 %temp; &#xD;&#xA;&#x9;mov.b64 &#x9;{%temp, %r6}, %fd15;&#xD;&#xA;&#x9;}&#xD;&#xA;&#x9;add.s32 &#x9;%r7, %r6, -1048576;&#xD;&#xA;&#x9;{&#xD;&#xA;&#x9;.reg .b32 %temp; &#xD;&#xA;&#x9;mov.b64 &#x9;{%r8, %temp}, %fd15;&#xD;&#xA;&#x9;}&#xD;&#xA;&#x9;mov.b64 &#x9;%fd17, {%r8, %r7};&#xD;&#xA;&#x9;mul.f64 &#x9;%fd18, %fd16, %fd15;&#xD;&#xA;&#x9;neg.f64 &#x9;%fd19, %fd18;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd20, %fd18, %fd19, %fd16;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd21, %fd20, %fd17, %fd18;&#xD;&#xA;&#x9;neg.f64 &#x9;%fd22, %fd21;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd23, %fd21, %fd22, %fd16;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd25, %fd15, %fd22, %fd14;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd26, %fd25, %fd17, %fd17;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd27, %fd23, %fd26, %fd21;&#xD;&#xA;&#x9;{&#xD;&#xA;&#x9;.reg .b32 %temp; &#xD;&#xA;&#x9;mov.b64 &#x9;{%temp, %r9}, %fd27;&#xD;&#xA;&#x9;}&#xD;&#xA;&#x9;add.s32 &#x9;%r10, %r9, 1048576;&#xD;&#xA;&#x9;{&#xD;&#xA;&#x9;.reg .b32 %temp; &#xD;&#xA;&#x9;mov.b64 &#x9;{%r11, %temp}, %fd27;&#xD;&#xA;&#x9;}&#xD;&#xA;&#x9;mov.b64 &#x9;%fd28, {%r11, %r10};&#xD;&#xA;&#x9;mov.f64 &#x9;%fd29, 0dBEBAC2FE66FAAC4B;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd30, 0d3EC715B371155F70;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd31, %fd30, %fd2, %fd29;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd32, 0d3ED9A9B88EFCD9B8;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd33, %fd31, %fd2, %fd32;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd34, 0d3EDD0F40A8A0C4C3;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd35, %fd33, %fd2, %fd34;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd36, 0d3EF46D4CFA9E0E1F;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd37, %fd35, %fd2, %fd36;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd38, 0d3F079C168D1E2422;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd39, %fd37, %fd2, %fd38;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd40, 0d3F1C9A88C3BCA540;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd41, %fd39, %fd2, %fd40;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd42, 0d3F31C4E64BD476DF;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd43, %fd41, %fd2, %fd42;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd44, 0d3F46E8BA60009C8F;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd45, %fd43, %fd2, %fd44;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd46, 0d3F5F1C71C62B05A2;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd47, %fd45, %fd2, %fd46;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd48, 0d3F76DB6DB6DC9F2C;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd49, %fd47, %fd2, %fd48;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd50, 0d3F9333333333329C;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd51, %fd49, %fd2, %fd50;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd52, 0d3FB5555555555555;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd53, %fd51, %fd2, %fd52;&#xD;&#xA;&#x9;mul.f64 &#x9;%fd54, %fd2, %fd53;&#xD;&#xA;&#x9;fma.rn.f64 &#x9;%fd97, %fd54, %fd28, %fd28;&#xD;&#xA;&#x9;bra.uni &#x9;$L__BB0_4;&#xD;&#xA;&#xD;&#xA;$L__BB0_10:&#xD;&#xA;&#x9;mov.f64 &#x9;%fd92, 0d3C91A62633145C07;&#xD;&#xA;&#x9;add.rn.f64 &#x9;%fd93, %fd9, %fd92;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd94, 0d3FF921FB54442D18;&#xD;&#xA;&#x9;add.rn.f64 &#x9;%fd97, %fd94, %fd93;&#xD;&#xA;&#x9;bra.uni &#x9;$L__BB0_11;&#xD;&#xA;&#xD;&#xA;$L__BB0_3:&#xD;&#xA;&#x9;mov.f64 &#x9;%fd55, 0d0000000000000000;&#xD;&#xA;&#x9;mul.rn.f64 &#x9;%fd97, %fd1, %fd55;&#xD;&#xA;&#xD;&#xA;$L__BB0_4:&#xD;&#xA;&#x9;setp.gt.s32 &#x9;%p3, %r2, -1;&#xD;&#xA;&#x9;@%p3 bra &#x9;$L__BB0_6;&#xD;&#xA;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd56, 0d7FF0000000000000;&#xD;&#xA;&#x9;mul.rn.f64 &#x9;%fd97, %fd97, %fd56;&#xD;&#xA;&#xD;&#xA;$L__BB0_6:&#xD;&#xA;&#x9;setp.gt.s32 &#x9;%p4, %r1, -1;&#xD;&#xA;&#x9;@%p4 bra &#x9;$L__BB0_11;&#xD;&#xA;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd57, 0dBCA1A62633145C07;&#xD;&#xA;&#x9;add.rn.f64 &#x9;%fd58, %fd97, %fd57;&#xD;&#xA;&#x9;neg.f64 &#x9;%fd59, %fd58;&#xD;&#xA;&#x9;mov.f64 &#x9;%fd60, 0d400921FB54442D18;&#xD;&#xA;&#x9;add.rn.f64 &#x9;%fd97, %fd60, %fd59;&#xD;&#xA;&#xD;&#xA;$L__BB0_11:&#xD;&#xA;&#x9;st.param.f64 &#x9;[func_retval0+0], %fd97;&#xD;&#xA;&#x9;ret;&#xD;&#xA;&#xD;&#xA;}" PtxDeclaration=".extern .func  (.param .b64 func_retval0) __ilgpu__nv_acos(&#xD;&#xA;&#x9;.param .b64 __ilgpu__nv_acos_param_0&#xD;&#xA;)&#xD;&#xA;;" />
   <Function Name="__nv_acosf" PtxModule="//&#xD;&#xA;// Generated by NVIDIA NVVM Compiler&#xD;&#xA;//&#xD;&#xA;// Compiler Build ID: CL-31968024&#xD;&#xA;// Cuda compilation tools, release 12.0, V12.0.76&#xD;&#xA;// Based on NVVM 7.0.1&#xD;&#xA;//&#xD;&#xA;&#xD;&#xA;.version 8.0&#xD;&#xA;.target sm_60&#xD;&#xA;.address_size 64&#xD;&#xA;&#xD;&#xA;&#x9;// .globl&#x9;__ilgpu__nv_acosf&#xD;&#xA;&#xD;&#xA;.visible .func  (.param .b32 func_retval0) __ilgpu__nv_acosf(&#xD;&#xA;&#x9;.param .b32 __ilgpu__nv_acosf_param_0&#xD;&#xA;)&#xD;&#xA;{&#xD;&#xA;&#x9;.reg .pred &#x9;%p&lt;3&gt;;&#xD;&#xA;&#x9;.reg .f32 &#x9;%f&lt;27&gt;;&#xD;&#xA;&#xD;&#xA;&#xD;&#xA;&#x9;ld.param.f32 &#x9;%f1, [__ilgpu__nv_acosf_param_0];&#xD;&#xA;&#x9;abs.f32 &#x9;%f2, %f1;&#xD;&#xA;&#x9;mov.f32 &#x9;%f3, 0f3F800000;&#xD;&#xA;&#x9;sub.f32 &#x9;%f4, %f3, %f2;&#xD;&#xA;&#x9;mul.f32 &#x9;%f5, %f4, 0f3F000000;&#xD;&#xA;&#x9;sqrt.rn.f32 &#x9;%f6, %f5;&#xD;&#xA;&#x9;setp.gt.f32 &#x9;%p1, %f2, 0f3F11EB85;&#xD;&#xA;&#x9;selp.f32 &#x9;%f7, %f6, %f2, %p1;&#xD;&#xA;&#x9;mul.f32 &#x9;%f8, %f7, %f7;&#xD;&#xA;&#x9;mov.f32 &#x9;%f9, 0f3C94D2E9;&#xD;&#xA;&#x9;mov.f32 &#x9;%f10, 0f3D53F941;&#xD;&#xA;&#x9;fma.rn.f32 &#x9;%f11, %f10, %f8, %f9;&#xD;&#xA;&#x9;mov.f32 &#x9;%f12, 0f3D3F841F;&#xD;&#xA;&#x9;fma.rn.f32 &#x9;%f13, %f11, %f8, %f12;&#xD;&#xA;&#x9;mov.f32 &#x9;%f14, 0f3D994929;&#xD;&#xA;&#x9;fma.rn.f32 &#x9;%f15, %f13, %f8, %f14;&#xD;&#xA;&#x9;mov.f32 &#x9;%f16, 0f3E2AAB94;&#xD;&#xA;&#x9;fma.rn.f32 &#x9;%f17, %f15, %f8, %f16;&#xD;&#xA;&#x9;mul.f32 &#x9;%f18, %f8, %f17;&#xD;&#xA;&#x9;fma.rn.f32 &#x9;%f19, %f18, %f7, %f7;&#xD;&#xA;&#x9;add.f32 &#x9;%f20, %f19, %f19;&#xD;&#xA;&#x9;mov.f32 &#x9;%f21, 0f3FC90FDB;&#xD;&#xA;&#x9;sub.f32 &#x9;%f22, %f21, %f19;&#xD;&#xA;&#x9;selp.f32 &#x9;%f23, %f20, %f22, %p1;&#xD;&#xA;&#x9;setp.lt.f32 &#x9;%p2, %f1, 0f00000000;&#xD;&#xA;&#x9;mov.f32 &#x9;%f24, 0f40490FDB;&#xD;&#xA;&#x9;sub.f32 &#x9;%f25, %f24, %f23;&#xD;&#xA;&#x9;selp.f32 &#x9;%f26, %f25, %f23, %p2;&#xD;&#xA;&#x9;st.param.f32 &#x9;[func_retval0+0], %f26;&#xD;&#xA;&#x9;ret;&#xD;&#xA;&#xD;&#xA;}" PtxDeclaration=".extern .func  (.param .b32 func_retval0) __ilgpu__nv_acosf(&#xD;&#xA;&#x9;.param .b32 __ilgpu__nv_acosf_param_0&#xD;&#xA;)&#xD;&#xA;;" />