Added support for OpenCL C versions.

Added backend support for OpenCL C 2.0.

Src/ILGPU/Backends/OpenCL/CLBackend.cs

@@ -24,7 +24,12 @@ namespace ILGPU.Backends.OpenCL
     /// <summary>
     /// Represents an OpenCL source backend.
     /// </summary>
-    public sealed partial class CLBackend : CodeGeneratorBackend<CLIntrinsic.Handler, CLCodeGenerator.GeneratorArgs, CLCodeGenerator, StringBuilder>
+    public sealed partial class CLBackend :
+        CodeGeneratorBackend<
+            CLIntrinsic.Handler,
+            CLCodeGenerator.GeneratorArgs,
+            CLCodeGenerator,
+            StringBuilder>
     {
         #region Nested Types
 
@@ -53,6 +58,15 @@ public IntrinsicSpecializerConfiguration(ContextFlags flags)
 
         #endregion
 
+        #region Static
+
+        /// <summary>
+        /// Represents the minimum OpenCL C version that is required.
+        /// </summary>
+        public static readonly CLCVersion MinimumVersion = new CLCVersion(2, 0);
+
+        #endregion
+
         #region Instance
 
         /// <summary>
@@ -156,7 +170,8 @@ public IntrinsicSpecializerConfiguration(ContextFlags flags)
             return new CLCompiledKernel(
                 Context,
                 entryPoint as SeparateViewEntryPoint,
-                clSource);
+                clSource,
+                MinimumVersion);
         }
 
         #endregion

Src/ILGPU/Backends/OpenCL/CLCVersion.cs

@@ -0,0 +1,132 @@
+// -----------------------------------------------------------------------------
+//                                    ILGPU
+//                     Copyright (c) 2016-2019 Marcel Koester
+//                                www.ilgpu.net
+//
+// File: CLCVersion.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.Text.RegularExpressions;
+
+namespace ILGPU.Backends.OpenCL
+{
+    /// <summary>
+    /// Represents an OpenCL C version.
+    /// </summary>
+    public readonly struct CLCVersion
+    {
+        #region Static
+
+        /// <summary>
+        /// The OpenCL C version 1.0.
+        /// </summary>
+        public static readonly CLCVersion CL10 = new CLCVersion(1, 0);
+
+        /// <summary>
+        /// The OpenCL C version 1.1.
+        /// </summary>
+        public static readonly CLCVersion CL11 = new CLCVersion(1, 2);
+
+        /// <summary>
+        /// The OpenCL C version 1.2.
+        /// </summary>
+        public static readonly CLCVersion CL12 = new CLCVersion(1, 2);
+
+        /// <summary>
+        /// The OpenCL C version 2.0.
+        /// </summary>
+        public static readonly CLCVersion CL20 = new CLCVersion(2, 0);
+
+        /// <summary>
+        /// The internal regex that is used to parse OpenCL C versions.
+        /// </summary>
+        private static readonly Regex VersionRegex = new Regex("\\s*(CL|OpenCL C)?\\s*([0-9]+).([0.9]+)");
+
+        /// <summary>
+        /// Tries to parse the given string expression into an OpenCL C version.
+        /// </summary>
+        /// <param name="expression">The expression to parse.</param>
+        /// <param name="version">The parsed version (if any).</param>
+        /// <returns>True, if the given expression could be parsed into an OpenCL C version.</returns>
+        public static bool TryParse(string expression, out CLCVersion version)
+        {
+            version = default;
+            var match = VersionRegex.Match(expression);
+            if (!match.Success)
+                return false;
+            version = new CLCVersion(
+                int.Parse(match.Groups[2].Value),
+                int.Parse(match.Groups[3].Value));
+            return true;
+        }
+
+        #endregion
+
+        #region Instance
+
+        /// <summary>
+        /// Constructs a new OpenCL C version.
+        /// </summary>
+        /// <param name="major">The major version.</param>
+        /// <param name="minor">The minor version.</param>
+        public CLCVersion(int major, int minor)
+        {
+            Major = major;
+            Minor = minor;
+        }
+
+        #endregion
+
+        #region Properties
+
+        /// <summary>
+        /// The major OpenCL C Version.
+        /// </summary>
+        public int Major { get; }
+
+        /// <summary>
+        /// The minor OpenCL C Version.
+        /// </summary>
+        public int Minor { get; }
+
+        #endregion
+
+        #region Object
+
+        /// <summary>
+        /// Returns the OpenCL C string representation that is compatible
+        /// with the OpenCL API.
+        /// </summary>
+        /// <returns>The string representation of this OpenCL C version.</returns>
+        public override string ToString() => $"CL{Major}.{Minor}";
+
+        #endregion
+
+        #region Operators
+
+        /// <summary>
+        /// Returns true if the first version is smaller than the second one.
+        /// </summary>
+        /// <param name="first">The first version.</param>
+        /// <param name="second">The second version.</param>
+        /// <returns>True, if the first version is smaller than the second one.</returns>
+        public static bool operator <(CLCVersion first, CLCVersion second) =>
+            first.Major < second.Major ||
+            first.Major == second.Major && first.Minor < second.Minor;
+
+        /// <summary>
+        /// Returns true if the first version is greater than the second one.
+        /// </summary>
+        /// <param name="first">The first version.</param>
+        /// <param name="second">The second version.</param>
+        /// <returns>True, if the first version is greater than the second one.</returns>
+        public static bool operator >(CLCVersion first, CLCVersion second) =>
+            first.Major > second.Major ||
+            first.Major == second.Major && first.Minor > second.Minor;
+
+        #endregion
+    }
+}

Src/ILGPU/Backends/OpenCL/CLCompiledKernel.cs

@@ -35,13 +35,16 @@ public sealed class CLCompiledKernel : CompiledKernel
         /// <param name="context">The associated context.</param>
         /// <param name="entryPoint">The entry point.</param>
         /// <param name="source">The source code.</param>
+        /// <param name="version">The OpenCL C version.</param>
         public CLCompiledKernel(
             Context context,
             SeparateViewEntryPoint entryPoint,
-            string source)
+            string source,
+            CLCVersion version)
             : base(context, entryPoint)
         {
             Source = source;
+            CVersion = version;
         }
 
         #endregion
@@ -53,6 +56,11 @@ public sealed class CLCompiledKernel : CompiledKernel
         /// </summary>
         public string Source { get; }
 
+        /// <summary>
+        /// Returns the used OpenCL C version.
+        /// </summary>
+        public CLCVersion CVersion { get; }
+
         /// <summary>
         /// Returns the internally used entry point.
         /// </summary>

Src/ILGPU/Backends/OpenCL/CLKernelFunctionGenerator.cs

@@ -73,7 +73,7 @@ public override void GenerateHeader(StringBuilder builder)
         public override void GenerateCode()
         {
             // Emit kernel declaration and parameter definitions
-            Builder.Append("__kernel void ");
+            Builder.Append("kernel void ");
             Builder.Append(CLCompiledKernel.EntryName);
             Builder.AppendLine("(");
 
@@ -85,7 +85,7 @@ public override void GenerateCode()
             {
                 // Emit a specialized pointer type
                 var elementType = viewParameters[i].ElementType;
-                Builder.Append("\t __global ");
+                Builder.Append("\t global ");
                 Builder.Append(TypeGenerator[elementType]);
                 Builder.Append(' ');
                 Builder.Append(CLInstructions.DereferenceOperation);

Src/ILGPU/Backends/OpenCL/CLTypeGenerator.cs

@@ -123,10 +123,9 @@ public void Visit(ViewType type)
             {
                 BeginStruct(type);
                 Builder.Append('\t');
-                Builder.Append(CLInstructions.GetAddressSpacePrefix(type.AddressSpace));
-                Builder.Append(' ');
                 Builder.Append(TypeGenerator[type.ElementType]);
-                Builder.Append(" *");
+                Builder.Append(CLInstructions.DereferenceOperation);
+                Builder.Append(' ');
                 Builder.Append(ViewPointerName);
                 Builder.AppendLine(";");
                 Builder.Append("\tint ");
@@ -363,8 +362,6 @@ private void DefinePointerType(TypeNode typeNode)
                 return;
 
             Builder.Append("typedef ");
-            Builder.Append(CLInstructions.GetAddressSpacePrefix(pointerType.AddressSpace));
-            Builder.Append(' ');
             Builder.Append(this[pointerType.ElementType]);
             Builder.Append(CLInstructions.DereferenceOperation);
             Builder.Append(' ');

Src/ILGPU/Runtime/OpenCL/CLAccelerator.cs

@@ -211,6 +211,14 @@ public CLAccelerator(Context context, CLAcceleratorId acceleratorId)
                 DeviceId,
                 CLDeviceInfoType.CL_DEVICE_TYPE);
 
+            // Determine the supported OpenCL C version
+            var clVersionString = CLAPI.GetDeviceInfo(
+                DeviceId,
+                CLDeviceInfoType.CL_DEVICE_OPENCL_C_VERSION);
+            if (!CLCVersion.TryParse(clVersionString, out CLCVersion version))
+                version = CLCVersion.CL10;
+            CVersion = version;
+
             // Max grid size
             int workItemDimensions = IntrinsicMath.Max(CLAPI.GetDeviceInfo<int>(
                 DeviceId,
@@ -308,6 +316,7 @@ private void InitVendorFeatures()
                 CLException.ThrowIfFailed(CLKernel.LoadKernel(
                     this,
                     DummyKernelSource,
+                    CVersion,
                     out IntPtr programPtr,
                     out IntPtr kernelPtr,
                     out var _));
@@ -345,6 +354,7 @@ private void InitSubGroupSupport(CLAcceleratorId acceleratorId)
             if (CLKernel.LoadKernel(
                 this,
                 DummySubGroupKernelSource,
+                CVersion,
                 out IntPtr programPtr,
                 out IntPtr kernelPtr,
                 out var _) == CLError.CL_SUCCESS)
@@ -423,6 +433,11 @@ private void InitSubGroupSupport(CLAcceleratorId acceleratorId)
         /// </summary>
         public int ClockRate { get; }
 
+        /// <summary>
+        /// Returns the supported OpenCL C version.
+        /// </summary>
+        public CLCVersion CVersion { get; }
+
         /// <summary>
         /// Returns true if this accelerator has sub-group support.
         /// </summary>

Src/ILGPU/Runtime/OpenCL/CLKernel.cs

@@ -31,13 +31,15 @@ public sealed class CLKernel : Kernel
         /// </summary>
         /// <param name="accelerator">The associated accelerator.</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>
         internal static CLError LoadKernel(
             CLAccelerator accelerator,
             string source,
+            CLCVersion version,
             out IntPtr programPtr,
             out IntPtr kernelPtr,
             out string errorLog)
@@ -51,8 +53,8 @@ public sealed class CLKernel : Kernel
             if (programError != CLError.CL_SUCCESS)
                 return programError;
 
-            // TODO: OpenCL compiler options
-            string options = string.Empty;
+            // Specify the OpenCL C version.
+            string options = "-cl-std=" + version.ToString();
 
             var buildError = CLAPI.BuildProgram(
                 programPtr,
@@ -112,6 +114,7 @@ public sealed class CLKernel : Kernel
             var errorCode = LoadKernel(
                 accelerator,
                 kernel.Source,
+                kernel.CVersion,
                 out programPtr,
                 out kernelPtr,
                 out var errorLog);
@@ -128,6 +131,7 @@ public sealed class CLKernel : Kernel
             CLException.ThrowIfFailed(LoadKernel(
                 accelerator,
                 kernel.Source,
+                kernel.CVersion,
                 out programPtr,
                 out kernelPtr,
                 out var _));