ClNvidia.cs

using System;
using System.Reflection;
using System.Runtime.InteropServices;
using OpenCL.Net;

namespace TinkerWorX.Simulation.MonoGame
{
    public enum D3dDeviceSource : uint
    {
        D3D11_DEVICE = 0x4019,
        D3D11_DXGI_ADAPTER = 0x401A,
    }

    public enum D3dDeviceSet : uint
    {
        PREFERRED_DEVICES_FOR_D3D11 = 0x401B,
        ALL_DEVICES_FOR_D3D11 = 0x401C,
    }

    public static class ContextExtensions
    {
        private static readonly FieldInfo ContextHandleField;

        static ContextExtensions()
        {
            ContextHandleField = typeof(Context).GetField("_handle", BindingFlags.Instance | BindingFlags.NonPublic);
        }

        public static IntPtr GetHandle(this Context context) => (IntPtr)ContextHandleField.GetValue(context);
    }

    public static class PlatformExtensions
    {
        private static readonly FieldInfo PlatformHandleField;

        static PlatformExtensions()
        {
            PlatformHandleField = typeof(Platform).GetField("_handle", BindingFlags.Instance | BindingFlags.NonPublic);
        }

        public static IntPtr GetHandle(this Platform platform) => (IntPtr)PlatformHandleField.GetValue(platform);
    }

    public static class CommandQueueExtensions
    {
        private static readonly FieldInfo CommandQueueHandleField;

        static CommandQueueExtensions()
        {
            CommandQueueHandleField = typeof(CommandQueue).GetField("_handle", BindingFlags.Instance | BindingFlags.NonPublic);
        }

        public static IntPtr GetHandle(this CommandQueue platform) => (IntPtr)CommandQueueHandleField.GetValue(platform);
    }

    public static class MemExtensions
    {
        private static readonly FieldInfo MemHandleField;

        static MemExtensions()
        {
            MemHandleField = typeof(Mem).GetField("_handle", BindingFlags.Instance | BindingFlags.NonPublic);
        }

        public static IntPtr GetHandle(this Mem platform) => (IntPtr)MemHandleField.GetValue(platform);
    }

    public static class MemHelper
    {
        private static readonly ConstructorInfo constructor;

        static MemHelper()
        {
            constructor = typeof(Mem).GetConstructor(BindingFlags.NonPublic | BindingFlags.Instance, null, new[] { typeof(IntPtr) }, null);
        }

        public static Mem Create(IntPtr handle) => (Mem)constructor.Invoke(new object[] { handle });
    }

    public static class ClNvidia
    {
        [UnmanagedFunctionPointer(CallingConvention.Cdecl)]
        private delegate ErrorCode clGetDeviceIDsFromD3D11NV_fn(
            IntPtr platform,
            D3dDeviceSource deviceSource,
            IntPtr d3dDevice,
            D3dDeviceSet deviceSet,
            uint numEntries,
            [Out] [MarshalAs(UnmanagedType.LPArray)] Device[] devices,
            out uint numDevices);

        private static clGetDeviceIDsFromD3D11NV_fn clGetDeviceIDsFromD3D11NV;

        [UnmanagedFunctionPointer(CallingConvention.Cdecl)]
        private delegate IntPtr clCreateFromD3D11Texture2DNV_fn(
              IntPtr context,
              MemFlags flags,
              IntPtr resource,
              uint subresource,
              [Out] [MarshalAs(UnmanagedType.I4)] out ErrorCode errcodeRet);

        private static readonly clCreateFromD3D11Texture2DNV_fn clCreateFromD3D11Texture2DNV;

        [UnmanagedFunctionPointer(CallingConvention.Cdecl)]
        private delegate ErrorCode clEnqueueAcquireD3D11ObjectsNV_fn(
           IntPtr commandQueue,
           uint num_objects,
           [In] [MarshalAs(UnmanagedType.LPArray, SizeParamIndex = 1)] IntPtr[] mem_objects,
           uint numEventsInWaitList,
           [In] [MarshalAs(UnmanagedType.LPArray, ArraySubType = UnmanagedType.SysUInt, SizeParamIndex = 3)] Event[] eventWaitList,
           [Out] [MarshalAs(UnmanagedType.Struct)] out Event e);

        private static readonly clEnqueueAcquireD3D11ObjectsNV_fn clEnqueueAcquireD3D11ObjectsNV;

        [UnmanagedFunctionPointer(CallingConvention.Cdecl)]
        private delegate ErrorCode clEnqueueReleaseD3D11ObjectsNV_fn(
            IntPtr commandQueue,
            uint num_objects,
            [In] [MarshalAs(UnmanagedType.LPArray, SizeParamIndex = 1)] IntPtr[] mem_objects,
            uint numEventsInWaitList,
            [In] [MarshalAs(UnmanagedType.LPArray, ArraySubType = UnmanagedType.SysUInt, SizeParamIndex = 3)] Event[] eventWaitList,
            [Out] [MarshalAs(UnmanagedType.Struct)] out Event e);

        private static readonly clEnqueueReleaseD3D11ObjectsNV_fn clEnqueueReleaseD3D11ObjectsNV;

        static ClNvidia()
        {
            clGetDeviceIDsFromD3D11NV = Marshal.GetDelegateForFunctionPointer<clGetDeviceIDsFromD3D11NV_fn>(Cl.GetExtensionFunctionAddress(nameof(clGetDeviceIDsFromD3D11NV)));
            clCreateFromD3D11Texture2DNV = Marshal.GetDelegateForFunctionPointer<clCreateFromD3D11Texture2DNV_fn>(Cl.GetExtensionFunctionAddress(nameof(clCreateFromD3D11Texture2DNV)));
            clEnqueueAcquireD3D11ObjectsNV = Marshal.GetDelegateForFunctionPointer<clEnqueueAcquireD3D11ObjectsNV_fn>(Cl.GetExtensionFunctionAddress(nameof(clEnqueueAcquireD3D11ObjectsNV)));
            clEnqueueReleaseD3D11ObjectsNV = Marshal.GetDelegateForFunctionPointer<clEnqueueReleaseD3D11ObjectsNV_fn>(Cl.GetExtensionFunctionAddress(nameof(clEnqueueReleaseD3D11ObjectsNV)));
        }

        public static ErrorCode GetDeviceIDs(Platform platform, D3dDeviceSource deviceSource, IntPtr d3dDevice, D3dDeviceSet deviceSet, uint numEntries, Device[] devices, out uint numDevices)
        {
            return clGetDeviceIDsFromD3D11NV(platform.GetHandle(), deviceSource, d3dDevice, deviceSet, numEntries, devices, out numDevices);
        }

        public static Device[] GetDeviceIDs(Platform platform, D3dDeviceSource deviceSource, IntPtr d3dDevice, D3dDeviceSet deviceSet, out ErrorCode error)
        {
            error = GetDeviceIDs(platform, deviceSource, d3dDevice, deviceSet, 0, null, out var deviceCount);
            if (error != ErrorCode.Success)
                return new Device[0];

            var deviceIds = new Device[deviceCount];
            error = GetDeviceIDs(platform, deviceSource, d3dDevice, deviceSet, deviceCount, deviceIds, out _);
            if (error != ErrorCode.Success)
                return new Device[0];

            return deviceIds;
        }

        public static Mem CreateFromD3D11Texture2D(Context context, MemFlags flags, IntPtr resource, uint subresource, out ErrorCode errcodeRet)
        {
            return MemHelper.Create(clCreateFromD3D11Texture2DNV.Invoke(context.GetHandle(), flags, resource, subresource, out errcodeRet));
        }

        public static ErrorCode EnqueueAcquireD3D11Objects(CommandQueue commandQueue, Mem image, uint numEventsInWaitList, Event[] eventWaitList, out Event e)
        {
            return clEnqueueAcquireD3D11ObjectsNV(commandQueue.GetHandle(), 1, new IntPtr[] { image.GetHandle() }, numEventsInWaitList, eventWaitList, out e);
        }

        public static ErrorCode EnqueueReleaseD3D11Objects(CommandQueue commandQueue, Mem image, uint numEventsInWaitList, Event[] eventWaitList, out Event e)
        {
            return clEnqueueReleaseD3D11ObjectsNV(commandQueue.GetHandle(), 1, new IntPtr[] { image.GetHandle() }, numEventsInWaitList, eventWaitList, out e);
        }
    }
}

OpenClDirectXHeatmap2d.cs

using Microsoft.Xna.Framework;
using OpenCL.Net;
using System;
using System.IO;
using System.Runtime.InteropServices;
using TinkerWorX.Simulation.MonoGame;

namespace TinkerWorX.Simulation.Thermal
{
    public class OpenClDirectXHeatmap2d : IHeatmap2d, IDisposable
    {
        private const int ARG_U0 = 0;
        private const int ARG_U1 = 1;
        private const int ARG_IMAGE = 2;
        private const int ARG_WIDTH = 3;
        private const int ARG_HEIGHT = 4;
        private const int ARG_DELTA_TIME = 5;

        private readonly float[] data;

        private readonly Program program;

        private readonly Kernel kernel;

        private readonly CommandQueue commandQueue;

        private bool isSynced;

        private int width;

        private int height;

        private IMem<float> u0;

        private IMem<float> u1;

        private Mem image;

        public OpenClDirectXHeatmap2d(int width, int height, IntPtr d3dImagePtr, Device device, Context context)
        {
            this.data = new float[width * height];
            this.program = Cl.CreateProgramWithSource(context, 1, new string[] { File.ReadAllText($"Thermal/{nameof(OpenClDirectXHeatmap2d)}.cl") }, null, out var error);
            if (error != ErrorCode.Success)
                throw new InvalidOperationException($"Cl.CreateProgramWithSource failed: {error}");

            error = Cl.BuildProgram(this.program, 1, new[] { device }, string.Empty, null, IntPtr.Zero);
            if (error != ErrorCode.Success)
            {
                if (error == ErrorCode.BuildProgramFailure)
                    throw new InvalidOperationException($"Cl.BuildProgram failed: {Cl.GetProgramBuildInfo(this.program, device, ProgramBuildInfo.Log, out error)}");
                throw new InvalidOperationException($"Cl.BuildProgram failed: {error}");
            }

            this.kernel = Cl.CreateKernel(this.program, "heatmap", out error);
            if (error != ErrorCode.Success)
                throw new InvalidOperationException($"Cl.CreateKernel failed: {error}");

            this.U0 = Cl.CreateBuffer(context, MemFlags.ReadWrite | MemFlags.CopyHostPtr, this.data, out error);
            if (error != ErrorCode.Success)
                throw new InvalidOperationException($"Cl.CreateBuffer({nameof(this.U0)}) failed: {error}");

            this.U1 = Cl.CreateBuffer(context, MemFlags.ReadWrite | MemFlags.CopyHostPtr, this.data, out error);
            if (error != ErrorCode.Success)
                throw new InvalidOperationException($"Cl.CreateBuffer({nameof(this.U1)}) failed: {error}");

            this.Image = ClNvidia.CreateFromD3D11Texture2D(context, MemFlags.WriteOnly, d3dImagePtr, 0, out error);
            if (error != ErrorCode.Success)
                throw new InvalidOperationException($"ClD3d11.CreateFromD3D11Texture2DKHR({nameof(this.U1)}) failed: {error}");

            this.Width = width;

            this.Height = height;

            this.DeltaTime = 1.00f;

            this.commandQueue = Cl.CreateCommandQueue(context, device, CommandQueueProperties.None, out error);
            if (error != ErrorCode.Success)
                throw new InvalidOperationException($"Cl.CreateCommandQueue failed: {error}");
        }

        public float this[int x, int y]
        {
            get => this.Get(x, y);
            set => this.Set(x, y, value);
        }

        private IMem<float> U0
        {
            get => this.u0;
            set
            {
                if (value == this.u0)
                    return;

                this.u0 = value;
                var error = Cl.SetKernelArg(this.kernel, ARG_U0, this.u0);
                if (error != ErrorCode.Success)
                    throw new InvalidOperationException($"Cl.SetKernelArg({ARG_U0}:{nameof(this.U0)}) failed: {error}");
            }
        }

        private IMem<float> U1
        {
            get => this.u1;
            set
            {
                if (value == this.u1)
                    return;

                this.u1 = value;
                var error = Cl.SetKernelArg(this.kernel, ARG_U1, this.u1);
                if (error != ErrorCode.Success)
                    throw new InvalidOperationException($"Cl.SetKernelArg({ARG_U0}:{nameof(this.U1)}) failed: {error}");
            }
        }

        private Mem Image
        {
            get => this.image;
            set
            {
                this.image = value;
                var error = Cl.SetKernelArg(this.kernel, ARG_IMAGE, this.image);
                if (error != ErrorCode.Success)
                    throw new InvalidOperationException($"Cl.SetKernelArg({ARG_IMAGE}:{nameof(this.Image)}) failed: {error}");
            }
        }

        public int Width
        {
            get => this.width;
            set
            {
                if (value == this.width)
                    return;

                this.width = value;
                var error = Cl.SetKernelArg(this.kernel, ARG_WIDTH, this.width);
                if (error != ErrorCode.Success)
                    throw new InvalidOperationException($"Cl.SetKernelArg({ARG_WIDTH}:{nameof(this.Width)}) failed: {error}");
            }
        }

        public int Height
        {
            get => this.height;
            set
            {
                if (value == this.height)
                    return;

                this.height = value;
                var error = Cl.SetKernelArg(this.kernel, ARG_HEIGHT, this.height);
                if (error != ErrorCode.Success)
                    throw new InvalidOperationException($"Cl.SetKernelArg({ARG_HEIGHT}:{nameof(this.Height)}) failed: {error}");
            }
        }

        private float DeltaTime
        {
            set
            {
                var error = Cl.SetKernelArg(this.kernel, ARG_DELTA_TIME, value);
                if (error != ErrorCode.Success)
                    throw new InvalidOperationException($"Cl.SetKernelArg({ARG_DELTA_TIME}:{nameof(this.DeltaTime)}) failed: {error}");
            }
        }

        public float Get(int x, int y) => this.data[x + this.Width * y];

        public void Set(int x, int y, float value)
        {
            this.data[x + this.Width * y] = value;
            this.isSynced = false;
        }

        private void SwapBuffers()
        {
            var temp = this.U0;
            this.U0 = this.U1;
            this.U1 = temp;
        }

        public void Update(float deltaTime)
        {
            this.SwapBuffers();

            ErrorCode clError;
            Event clEvent;
            this.DeltaTime = deltaTime;

            // check if anything has changed
            if (!this.isSynced)
            {
                // write our new data
                clError = Cl.EnqueueWriteBuffer(this.commandQueue, this.u0, Bool.False, 0, this.data.Length, this.data, 0, null, out clEvent);
                clEvent.Dispose();
                if (clError != ErrorCode.Success)
                    throw new InvalidOperationException($"Cl.EnqueueWriteBuffer failed: {clError}");
                this.isSynced = true;
            }

            // acquires a lock on the texture, ensuring it's not in use by anything else
            clError = ClNvidia.EnqueueAcquireD3D11Objects(this.commandQueue, this.image, 0, null, out clEvent);
            clEvent.Dispose();
            if (clError != ErrorCode.Success)
                throw new InvalidOperationException($"Cl.EnqueueAcquireD3D11Objects failed: {clError}");

            // run the kernel
            clError = Cl.EnqueueNDRangeKernel(this.commandQueue, this.kernel, 2, new[] { (IntPtr)1, (IntPtr)1 }, new[] { (IntPtr)this.Width - 2, (IntPtr)this.Height - 2 }, null, 0, null, out clEvent);
            clEvent.Dispose();
            if (clError != ErrorCode.Success)
                throw new InvalidOperationException($"Cl.EnqueueNDRangeKernel failed: {clError}");

            // releases the lock on the texture, letting everything else use it again
            clError = ClNvidia.EnqueueReleaseD3D11Objects(this.commandQueue, this.image, 0, null, out clEvent);
            clEvent.Dispose();
            if (clError != ErrorCode.Success)
                throw new InvalidOperationException($"Cl.EnqueueAcquireD3D11Objects failed: {clError}");

            // read back our output data
            clError = Cl.EnqueueReadBuffer(this.commandQueue, this.u1, Bool.False, 0, this.data.Length, this.data, 0, null, out clEvent);
            clEvent.Dispose();
            if (clError != ErrorCode.Success)
                throw new InvalidOperationException($"Cl.EnqueueReadBuffer failed: {clError}");

            // wait for things to finish
            clError = Cl.Finish(this.commandQueue);
            if (clError != ErrorCode.Success)
                throw new InvalidOperationException($"Cl.Finish failed: {clError}");

            Cl.Flush(this.commandQueue);
        }

        public void Dispose()
        {
            this.program.Dispose();
            this.kernel.Dispose();
            this.u0?.Dispose();
            this.u1?.Dispose();
            this.commandQueue.Dispose();
        }
    }
}

OpenClHeatmap2d.cl

__kernel void heatmap(global float* u0, global float* u1, int width, int height, float dt)
{
    // find where we are in the kernel
    int x = get_global_id(0); // [1 ..  width - 1]
    int y = get_global_id(1); // [1 .. height - 1]
    int id = x + width * y;

    // look up the different values needed (row major 1d array)
    float t = u0[id - width]; // top    (x, y - 1)
    float l = u0[id - 1];     // left   (x - 1, y)
    float v = u0[id];         // center (x, y)
    float r = u0[id + 1];     // right  (x + 1, y)
    float b = u0[id + width]; // bottom (x, y + 1)

    // do some math
    float a = 4.00;

    float dx2 = l - 2 * v + r;
    float dy2 = t - 2 * v + b;

    u1[id] = v + a * dt * (dx2 + dy2);
}