1.Hybridizer HOWTO — Hello World

Our hello world is the addition of two vector of elements.

The method gets an EntryPoint attribute, which indicates that the function will be called from host code and will be run on device code. We can use the Parallel.For construct, which is natively proposed by .Net:


[EntryPoint]
public static void VectorAdd(double[] a, double[] b, int N)
{
    Parallel.For(0, N, (i) => { a[i] += b[i] });
}

We can also use explicit work distribution, which is done using the constructs proposed by CUDA: threadIdx/blockDim, blockIdx/gridDim. This is customizable and names can be changed, but the concept is similar:


[EntryPoint]
public static void VectorAdd(double[] a, double[] b, int N)
    for (int k = threadIdx.x + blockDim.x * blockIdx.x ;
        k < count ; k += blockDim.x * gridDim.x)
    {
        a[k] += b[k];
    }
}

Explicit work distribution can be used (for example) to distribute work among a 2D-grid.

2.Hybridizer HOW TO — Intrinsics

Using existing functions, whether provided by built-in C libraries or defined later-on in native C is done using IntrinsicFunction attribute. When generating the source code, the function call is replaced by the IntrinsicFunction name, and the contents of that function is ignored.


[IntrinsicFunction("erf")]
public double Erf(double x)
{
    ...
}

[EntryPoint("TestErfc")]
public void errorFunction(int count, double[] a, double[] b)
{
    for (int k = threadIdx.x + blockDim.x * blockIdx.x;
        k < count ; k += blockDim.x * gridDim.x)
    {
        b[k] = Erf (a[k]);
    }
}

3.Hybridizer HOWTO — Libraries Integration

It is also possible to use/integrate existing libraries for which device functions are defined, hence extending the concept of intrinsic functions to intrinsic types.

Note that in that case, the functions do not need an implementation if no behavior is expected in plain C#.


[IntrinsicType("curandStateMRG32k3a_t")]
[IntrinsicIncludeCUDA("curand_kernel.h")]
[StructLayout(LayoutKind.Sequential)]
public unsafe struct curandStateMRG32k3a_t
{
    public fixed double s1[3];
    public fixed double s2[3];
    public int boxmuller_flag;
    public int boxmuller_flag_double;
    public float boxmuller_extra;
    public double boxmuller_extra_double;
    [IntrinsicFunction("curand_init")]
    public static void curand_init(ulong seed,
        ulong subsequence, ulong offset,
        out curandStateMRG32k3a_t state)
    { throw new NotImplementedException(); }
    [IntrinsicFunction("curand")] public uint curand()
    { throw new NotImplementedException(); }
    [IntrinsicFunction("curand_log_normal")]
    public float curand_log_normal(float mean, float stdev)
    { throw new NotImplementedException(); }

} 

4.Hybridizer HOWTO — Printf and Builtins

The concept of intrinsics allows extensions and more control on how the code is generated. We also extend this concept for existing methods for which we do not have control on the code. The equivalent of the attribute is described in a builtin file. This allows the use of Console.Out.Write and generate printf. It also allows us to use System.Math.Exp which would be replaced by exp from cmath.


[EntryPoint("TestPrintf")]
public void testPrintf()
{
    Console.Out.WriteLine(
            "Comment from Thread {0} Block {1}",
        threadIdx.x, blockIdx.x);
}

[EntryPoint("TestExp")]
public void TestExp()
{
    exp = System.Math.Exp(1.0);
}