diff --git a/Src/ILGPU.Tests/KernelEntryPoints.cs b/Src/ILGPU.Tests/KernelEntryPoints.cs
index 68ed0148f..296c6085a 100644
--- a/Src/ILGPU.Tests/KernelEntryPoints.cs
+++ b/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);
+        }
     }
 }
diff --git a/Src/ILGPU/Backends/CompiledKernel.cs b/Src/ILGPU/Backends/CompiledKernel.cs
index 3f12edf9c..5c950fe81 100644
--- a/Src/ILGPU/Backends/CompiledKernel.cs
+++ b/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>
diff --git a/Src/ILGPU/Backends/EntryPoints/EntryPoint.cs b/Src/ILGPU/Backends/EntryPoints/EntryPoint.cs
index 7c82bac0c..c7c9c31e1 100644
--- a/Src/ILGPU/Backends/EntryPoints/EntryPoint.cs
+++ b/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>
diff --git a/Src/ILGPU/Backends/EntryPoints/EntryPointDescription.cs b/Src/ILGPU/Backends/EntryPoints/EntryPointDescription.cs
index 040142c36..ab405b2af 100644
--- a/Src/ILGPU/Backends/EntryPoints/EntryPointDescription.cs
+++ b/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>
diff --git a/Src/ILGPU/Backends/EntryPoints/KernelNameAttribute.cs b/Src/ILGPU/Backends/EntryPoints/KernelNameAttribute.cs
new file mode 100644
index 000000000..0ebd02bb2
--- /dev/null
+++ b/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
+    }
+}
diff --git a/Src/ILGPU/Backends/OpenCL/CLCompiledKernel.cs b/Src/ILGPU/Backends/OpenCL/CLCompiledKernel.cs
index fbd8087d5..5c2f05b5c 100644
--- a/Src/ILGPU/Backends/OpenCL/CLCompiledKernel.cs
+++ b/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
diff --git a/Src/ILGPU/Backends/OpenCL/CLKernelFunctionGenerator.cs b/Src/ILGPU/Backends/OpenCL/CLKernelFunctionGenerator.cs
index b11cb3b53..1a28de94f 100644
--- a/Src/ILGPU/Backends/OpenCL/CLKernelFunctionGenerator.cs
+++ b/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
diff --git a/Src/ILGPU/Backends/PTX/PTXCodeGenerator.cs b/Src/ILGPU/Backends/PTX/PTXCodeGenerator.cs
index 8cb6c8ed9..8ee18e7b6 100644
--- a/Src/ILGPU/Backends/PTX/PTXCodeGenerator.cs
+++ b/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.
diff --git a/Src/ILGPU/Backends/PTX/PTXCompiledKernel.cs b/Src/ILGPU/Backends/PTX/PTXCompiledKernel.cs
index ee0b8a754..f81122584 100644
--- a/Src/ILGPU/Backends/PTX/PTXCompiledKernel.cs
+++ b/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
diff --git a/Src/ILGPU/Backends/PTX/PTXKernelFunctionGenerator.cs b/Src/ILGPU/Backends/PTX/PTXKernelFunctionGenerator.cs
index 236f2dfc4..eb44630a5 100644
--- a/Src/ILGPU/Backends/PTX/PTXKernelFunctionGenerator.cs
+++ b/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);
diff --git a/Src/ILGPU/Runtime/Cuda/CudaKernel.cs b/Src/ILGPU/Runtime/Cuda/CudaKernel.cs
index d6c437ee0..4b445a304 100644
--- a/Src/ILGPU/Runtime/Cuda/CudaKernel.cs
+++ b/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
diff --git a/Src/ILGPU/Runtime/OpenCL/CLAccelerator.cs b/Src/ILGPU/Runtime/OpenCL/CLAccelerator.cs
index 9a487716e..168ec2c84 100644
--- a/Src/ILGPU/Runtime/OpenCL/CLAccelerator.cs
+++ b/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,
diff --git a/Src/ILGPU/Runtime/OpenCL/CLKernel.cs b/Src/ILGPU/Runtime/OpenCL/CLKernel.cs
index d102a3076..0d3c55c35 100644
--- a/Src/ILGPU/Runtime/OpenCL/CLKernel.cs
+++ b/Src/ILGPU/Runtime/OpenCL/CLKernel.cs
@@ -12,7 +12,6 @@
 using ILGPU.Backends.OpenCL;
 using System;
 using System.Diagnostics;
-using System.Diagnostics.CodeAnalysis;
 using System.Reflection;
 using static ILGPU.Runtime.OpenCL.CLAPI;
 
@@ -37,12 +36,43 @@ public sealed class CLKernel : Kernel
         /// <returns>
         /// True, if the program and the kernel could be loaded successfully.
         /// </returns>
+        [Obsolete("Use LoadKernel with an explicit entry point name instead.")]
         public static CLError LoadKernel(
             CLAccelerator accelerator,
             string source,
             CLCVersion version,
             out IntPtr programPtr,
             out IntPtr kernelPtr,
+            out string errorLog) =>
+            LoadKernel(
+                accelerator,
+                CLCompiledKernel.EntryName,
+                source,
+                version,
+                out programPtr,
+                out kernelPtr,
+                out errorLog);
+
+        /// <summary>
+        /// Loads the given OpenCL kernel.
+        /// </summary>
+        /// <param name="accelerator">The associated accelerator.</param>
+        /// <param name="name">The name of the entry-point function.</param>
+        /// <param name="source">The OpenCL source code.</param>
+        /// <param name="version">The OpenCL C version.</param>
+        /// <param name="programPtr">The created program pointer.</param>
+        /// <param name="kernelPtr">The created kernel pointer.</param>
+        /// <param name="errorLog">The error log (if any).</param>
+        /// <returns>
+        /// True, if the program and the kernel could be loaded successfully.
+        /// </returns>
+        public static CLError LoadKernel(
+            CLAccelerator accelerator,
+            string name,
+            string source,
+            CLCVersion version,
+            out IntPtr programPtr,
+            out IntPtr kernelPtr,
             out string errorLog)
         {
             errorLog = null;
@@ -77,7 +107,7 @@ public static CLError LoadKernel(
 
             return CurrentAPI.CreateKernel(
                 programPtr,
-                CLCompiledKernel.EntryName,
+                name,
                 out kernelPtr);
         }
 
@@ -133,20 +163,15 @@ public static unsafe byte[] LoadBinaryRepresentation(IntPtr program)
         /// <param name="accelerator">The associated accelerator.</param>
         /// <param name="kernel">The source kernel.</param>
         /// <param name="launcher">The launcher method for the given kernel.</param>
-        [SuppressMessage(
-            "Microsoft.Design",
-            "CA1062:Validate arguments of public methods",
-            MessageId = "0",
-            Justification = "Will be verified in the constructor of the base class")]
         public CLKernel(
             CLAccelerator accelerator,
             CLCompiledKernel kernel,
             MethodInfo launcher)
             : base(accelerator, kernel, launcher)
         {
-#if DEBUG
             var errorCode = LoadKernel(
                 accelerator,
+                kernel.Name,
                 kernel.Source,
                 kernel.CVersion,
                 out programPtr,
@@ -154,23 +179,13 @@ public CLKernel(
                 out var errorLog);
             if (errorCode != CLError.CL_SUCCESS)
             {
-                Debug.WriteLine("Kernel loading failed:");
+                Trace.WriteLine("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);
             }
             CLException.ThrowIfFailed(errorCode);
-#else
-            CLException.ThrowIfFailed(LoadKernel(
-                accelerator,
-                kernel.Source,
-                kernel.CVersion,
-                out programPtr,
-                out kernelPtr,
-                out var _));
-#endif
-
         }
 
         #endregion