Added support for custom kernel names.

Src/ILGPU.Tests/KernelEntryPoints.cs

@@ -1,4 +1,5 @@
-using System;
+using ILGPU.Backends.EntryPoints;
+using System;
 using System.Diagnostics.CodeAnalysis;
 using System.Linq;
 using Xunit;
@@ -535,5 +536,24 @@ public void StaticFieldCapturingLambdaIndex3EntryPoint(int length)
                 Execute(kernel.Method, extent, buffer.View, extent));
             Assert.IsType<NotSupportedException>(e.InnerException);
         }
+
+        [KernelName("My @ CustomKernel.Name12345 [1211]")]
+        internal static void NamedEntryPointKernel(Index1 index, ArrayView<int> output)
+        {
+            output[index] = index;
+        }
+
+        [Fact]
+        [KernelMethod(nameof(NamedEntryPointKernel))]
+        public void NamedEntryPoint()
+        {
+            const int Length = 32;
+
+            using var buffer = Accelerator.Allocate<int>(Length);
+            Execute(buffer.Length, buffer.View);
+
+            var expected = Enumerable.Range(0, Length).ToArray();
+            Verify(buffer, expected);
+        }
     }
 }

Src/ILGPU/Backends/CompiledKernel.cs

@@ -204,6 +204,11 @@ protected CompiledKernel(
         /// </summary>
         public MethodInfo SourceMethod => EntryPoint.MethodInfo;
 
+        /// <summary>
+        /// Returns the associated kernel function name.
+        /// </summary>
+        public string Name => EntryPoint.Name;
+
         /// <summary>
         /// Returns the index type of the entry point.
         /// </summary>

Src/ILGPU/Backends/EntryPoints/EntryPoint.cs

@@ -53,6 +53,11 @@ public EntryPoint(
         /// </summary>
         public EntryPointDescription Description { get; }
 
+        /// <summary>
+        /// Returns the associated kernel function name.
+        /// </summary>
+        public string Name => Description.Name;
+
         /// <summary>
         /// Returns the associated method info.
         /// </summary>

Src/ILGPU/Backends/EntryPoints/EntryPointDescription.cs

@@ -101,6 +101,7 @@ internal EntryPointDescription(
                 parameterTypes.Add(type);
             }
             Parameters = new ParameterCollection(parameterTypes.MoveToImmutable());
+
             Validate();
         }
 
@@ -113,6 +114,12 @@ internal EntryPointDescription(
         /// </summary>
         public MethodInfo MethodSource { get; }
 
+        /// <summary>
+        /// Returns the name of the underlying entry point to be used in the scope of
+        /// loaded runtime <see cref="Kernel"/> instances.
+        /// </summary>
+        public readonly string Name => KernelNameAttribute.GetKernelName(MethodSource);
+
         /// <summary>
         /// Returns the associated index type.
         /// </summary>

Src/ILGPU/Backends/EntryPoints/KernelNameAttribute.cs

@@ -0,0 +1,95 @@
+// ---------------------------------------------------------------------------------------
+//                                        ILGPU
+//                        Copyright (c) 2016-2020 Marcel Koester
+//                                    www.ilgpu.net
+//
+// File: KernelNameAttribute.cs
+//
+// This file is part of ILGPU and is distributed under the University of Illinois Open
+// Source License. See LICENSE.txt for details
+// ---------------------------------------------------------------------------------------
+
+using System;
+using System.Reflection;
+
+namespace ILGPU.Backends.EntryPoints
+{
+    /// <summary>
+    /// Specifies a custom kernel name used in OpenCL or PTX kernels.
+    /// </summary>
+    /// <remarks>
+    /// Kernel names have to consist of ASCII characters only.
+    /// </remarks>
+    [AttributeUsage(AttributeTargets.Method, AllowMultiple = false)]
+    public sealed class KernelNameAttribute : Attribute
+    {
+        #region Constants
+
+        /// <summary>
+        /// The internally used kernel prefix to avoid clashes with other utility/local
+        /// functions in the finally emitted assembly.
+        /// </summary>
+        private const string KernelNamePrefix = "Kernel_";
+
+        #endregion
+
+        #region Static
+
+        /// <summary>
+        /// Gets the kernel name for the given entry point function.
+        /// </summary>
+        /// <param name="methodInfo">The entry point function.</param>
+        /// <returns>The kernel name.</returns>
+        public static string GetKernelName(MethodInfo methodInfo)
+        {
+            if (methodInfo is null)
+                throw new ArgumentNullException(nameof(methodInfo));
+            var attribute = methodInfo.GetCustomAttribute<KernelNameAttribute>();
+            var kernelName = GetCompatibleName(attribute?.KernelName ?? methodInfo.Name);
+            return KernelNamePrefix + kernelName;
+        }
+
+        /// <summary>
+        /// Returns a compatible function name for all runtime backends.
+        /// </summary>
+        /// <param name="name">The source name.</param>
+        internal static string GetCompatibleName(string name)
+        {
+            var chars = name.ToCharArray();
+            for (int i = 0, e = chars.Length; i < e; ++i)
+            {
+                ref var charValue = ref chars[i];
+                // Map to ASCII and letter/digit characters only
+                if (charValue >= 128 || !char.IsLetterOrDigit(charValue))
+                    charValue = '_';
+            }
+            return new string(chars);
+        }
+
+        #endregion
+
+        #region Instance
+
+        /// <summary>
+        /// Constructs a new kernel name attribute.
+        /// </summary>
+        /// <param name="kernelName">The kernel name to use.</param>
+        public KernelNameAttribute(string kernelName)
+        {
+            if (string.IsNullOrWhiteSpace(kernelName))
+                throw new ArgumentNullException(nameof(kernelName));
+            KernelName = GetCompatibleName(kernelName);
+        }
+
+        #endregion
+
+        #region Properties
+
+        /// <summary>
+        /// Returns the kernel name to use.
+        /// </summary>
+        public string KernelName { get; }
+
+        #endregion
+    }
+}

Src/ILGPU/Backends/OpenCL/CLCompiledKernel.cs

@@ -10,6 +10,7 @@
 // ---------------------------------------------------------------------------------------
 
 using ILGPU.Backends.EntryPoints;
+using System;
 
 namespace ILGPU.Backends.OpenCL
 {
@@ -23,6 +24,7 @@ public sealed class CLCompiledKernel : CompiledKernel
         /// <summary>
         /// The entry name of the kernel function.
         /// </summary>
+        [Obsolete("Use CompiledKernel.Name instead")]
         public const string EntryName = "ILGPUKernel";
 
         #endregion

Src/ILGPU/Backends/OpenCL/CLKernelFunctionGenerator.cs

@@ -206,7 +206,7 @@ public override void GenerateCode()
         {
             // Emit kernel declaration and parameter definitions
             Builder.Append("kernel void ");
-            Builder.Append(CLCompiledKernel.EntryName);
+            Builder.Append(EntryPoint.Name);
             Builder.AppendLine("(");
 
             // Initialize view information

Src/ILGPU/Backends/PTX/PTXCodeGenerator.cs

@@ -281,17 +281,8 @@ private static BasicValueType ResolveIOType(
         /// <param name="name">The source name.</param>
         /// <param name="nodeId">The source node id.</param>
         /// <returns>The resolved PTX name.</returns>
-        private static string GetCompatibleName(string name, NodeId nodeId)
-        {
-            var chars = name.ToCharArray();
-            for (int i = 0, e = chars.Length; i < e; ++i)
-            {
-                ref var charValue = ref chars[i];
-                if (!char.IsLetterOrDigit(charValue))
-                    charValue = '_';
-            }
-            return new string(chars) + nodeId.ToString();
-        }
+        private static string GetCompatibleName(string name, NodeId nodeId) =>
+            KernelNameAttribute.GetCompatibleName(name) + nodeId.ToString();
 
         /// <summary>
         /// Returns the PTX function name for the given function.

Src/ILGPU/Backends/PTX/PTXCompiledKernel.cs

@@ -10,6 +10,7 @@
 // ---------------------------------------------------------------------------------------
 
 using ILGPU.Backends.EntryPoints;
+using System;
 
 namespace ILGPU.Backends.PTX
 {
@@ -23,6 +24,7 @@ public sealed class PTXCompiledKernel : CompiledKernel
         /// <summary>
         /// The entry name of the kernel function.
         /// </summary>
+        [Obsolete("Use CompiledKernel.Name instead")]
         public const string EntryName = "ILGPUKernel";
 
         #endregion

Src/ILGPU/Backends/PTX/PTXKernelFunctionGenerator.cs

@@ -147,7 +147,7 @@ public override void GenerateCode()
         {
             Builder.AppendLine();
             Builder.Append(".visible .entry ");
-            Builder.Append(PTXCompiledKernel.EntryName);
+            Builder.Append(EntryPoint.Name);
             Builder.AppendLine("(");
 
             var parameterLogic = new KernelParameterSetupLogic(EntryPoint, this);

Src/ILGPU/Runtime/Cuda/CudaKernel.cs

@@ -48,31 +48,25 @@ internal CudaKernel(
             MethodInfo launcher)
             : base(accelerator, kernel, launcher)
         {
-#if DEBUG
             var kernelLoaded = CurrentAPI.LoadModule(
                 out modulePtr,
                 kernel.PTXAssembly,
                 out string errorLog);
             if (kernelLoaded != CudaError.CUDA_SUCCESS)
             {
-                Debug.WriteLine("Kernel loading failed:");
+                Trace.WriteLine("PTX Kernel loading failed:");
                 if (string.IsNullOrWhiteSpace(errorLog))
-                    Debug.WriteLine(">> No error information available");
+                    Trace.WriteLine(">> No error information available");
                 else
-                    Debug.WriteLine(errorLog);
+                    Trace.WriteLine(errorLog);
             }
             CudaException.ThrowIfFailed(kernelLoaded);
-#else
-            CudaException.ThrowIfFailed(
-                CurrentAPI.LoadModule(
-                    out modulePtr,
-                    kernel.PTXAssembly));
-#endif
+
             CudaException.ThrowIfFailed(
                 CurrentAPI.GetModuleFunction(
                     out functionPtr,
                     modulePtr,
-                    PTXCompiledKernel.EntryName));
+                    kernel.Name));
         }
 
         #endregion

Src/ILGPU/Runtime/OpenCL/CLAccelerator.cs

@@ -89,11 +89,16 @@ public sealed class CLAccelerator : KernelAccelerator<CLCompiledKernel, CLKernel
                 nameof(CLException.ThrowIfFailed),
                 BindingFlags.Public | BindingFlags.Static);
 
+        /// <summary>
+        /// Specifies the kernel entry point name for the following dummy kernels.
+        /// </summary>
+        private const string DummyKernelName = "ILGPUTestKernel";
+
         /// <summary>
         /// The first dummy kernel that is compiled during accelerator initialization.
         /// </summary>
         private const string DummyKernelSource =
-            "__kernel void " + CLCompiledKernel.EntryName + "(\n" +
+            "__kernel void " + DummyKernelName + "(\n" +
             "   __global const int *a,\n" +
             "   __global const int *b,\n" +
             "   __global int *c) { \n" +
@@ -104,7 +109,7 @@ public sealed class CLAccelerator : KernelAccelerator<CLCompiledKernel, CLKernel
         /// The second dummy kernel that is compiled during accelerator initialization.
         /// </summary>
         private const string DummySubGroupKernelSource =
-            "__kernel void " + CLCompiledKernel.EntryName + "(\n" +
+            "__kernel void " + DummyKernelName + "(\n" +
             "   __global int *a," +
             "   const int n) { \n" +
             "   size_t i = get_global_id(0);\n" +
@@ -357,6 +362,7 @@ private void InitVendorFeatures()
                 // Compile dummy kernel to resolve additional information
                 CLException.ThrowIfFailed(CLKernel.LoadKernel(
                     this,
+                    DummyKernelName,
                     DummyKernelSource,
                     CVersion,
                     out IntPtr programPtr,
@@ -395,6 +401,7 @@ private void InitSubGroupSupport(CLAcceleratorId acceleratorId)
             // Verify support using a simple kernel
             if (CLKernel.LoadKernel(
                 this,
+                DummyKernelName,
                 DummySubGroupKernelSource,
                 CVersion,
                 out IntPtr programPtr,