internal static void ArrayViewVectorizedIOKernel <T, T2>(
Index1D index,
ArrayView1D <T, Stride1D.Dense> source,
ArrayView1D <T, Stride1D.Dense> target)
where T : unmanaged
where T2 : unmanaged
{
// Use compile-time known offsets to test the internal alignment rules
var nonVectorAlignedSource = source.SubView(1, 2);
var nonVectorAlignedCastedSource = nonVectorAlignedSource.Cast <T, T2>();
var nonVectorAlignedTarget = target.SubView(1, 2);
var nonVectorAlignedTargetCasted = nonVectorAlignedTarget.Cast <T, T2>();
// Load from source and write to target
T2 data = nonVectorAlignedCastedSource[index];
nonVectorAlignedTargetCasted[index] = data;
// Perform the same operations with compile-time known offsets
var vectorAlignedSource = source.SubView(2, 2);
var vectorAlignedCastedSource = vectorAlignedSource.Cast <T, T2>();
var vectorAlignedTarget = target.SubView(2, 2);
var vectorAlignedTargetCasted = vectorAlignedTarget.Cast <T, T2>();
// Load from source and write to target
T2 data2 = vectorAlignedCastedSource[index];
vectorAlignedTargetCasted[index] = data2;
}
/// <summary>
/// Implicitly-grouped kernels receive an index type (first parameter) of type:
/// <see cref="Index"/>, <see cref="Index2"/> or <see cref="Index3"/>.
/// These kernel types hide the underlying blocking/grouping semantics of a GPU
/// and allow convenient kernel programming without having take grouping details into account.
/// The block or group size can be defined while loading a kernel via:
/// - LoadImplicitlyGroupedKernel
/// - LoadAutoGroupedKernel.
///
/// Note that you must not use warp-shuffle functionality within implicitly-grouped
/// kernels since not all lanes of a warp are guaranteed to participate in the warp shuffle.
/// </summary>
/// <param name="index">The current thread index.</param>
/// <param name="dataView">The view pointing to our memory buffer.</param>
/// <param name="constant">A nice uniform constant.</param>
static void MyKernel(
Index1D index, // The global thread index (1D in this case)
ArrayView <int> dataView, // A view to a chunk of memory (1D in this case)
int constant) // A uniform constant
{
dataView[index] = index + constant;
}
internal static void DebugAssertKernel(
Index1D index,
ArrayView1D <int, Stride1D.Dense> data)
{
Debug.Assert(data[index] >= 0);
Trace.Assert(data[index] >= 0);
}
internal static void NestedStructureCallOutKernel(
Index1D index,
ArrayView1D <int, Stride1D.Dense> data)
{
GetStructureValue(out Parent value);
data[index] = value.Second.Value;
}
/// <summary>
/// ILGPU kernel for Mandelbrot set.
/// </summary>
/// <param name="index"></param>
/// <param name="width"></param>
/// <param name="height"></param>
/// <param name="max_iterations"></param>
/// <param name="output"></param>
static void MandelbrotKernel(
Index1D index,
int width, int height, int max_iterations,
ArrayView <int> output)
{
float h_a = -2.0f;
float h_b = 1.0f;
float v_a = -1.0f;
float v_b = 1.0f;
if (index >= output.Length)
{
return;
}
int img_x = index % width;
int img_y = index / width;
float x0 = h_a + img_x * (h_b - h_a) / width;
float y0 = v_a + img_y * (v_b - v_a) / height;
float x = 0.0f;
float y = 0.0f;
int iteration = 0;
while ((x * x + y * y < 2 * 2) && (iteration < max_iterations))
{
float xtemp = x * x - y * y + x0;
y = 2 * x * y + y0;
x = xtemp;
iteration += 1;
}
output[index] = iteration;
}
internal static void MultiDimArraySimpleKernel <T, TArraySize>(
Index1D index,
ArrayView <T> dataX,
ArrayView <T> dataY,
ArrayView <T> dataZ,
ArrayView <T> dataW,
T c,
int localIndex)
where T : unmanaged
where TArraySize : unmanaged, ILength
{
TArraySize arraySize = default;
var array = new T[
arraySize.Length,
arraySize.Length,
arraySize.Length,
arraySize.Length];
for (int i = 0; i < array.GetLength(2); ++i)
{
array[i, i, i, i] = c;
}
dataX[index] = array[0, 0, 0, localIndex];
dataY[index] = array[0, 0, localIndex, 0];
dataZ[index] = array[0, localIndex, 0, 0];
dataW[index] = array[localIndex, 0, 0, 0];
}
internal static void NestedBreakContinueKernel(
Index1D index,
ArrayView1D <int, Stride1D.Dense> data,
int counter,
int counter2,
int counter3)
{
int accumulate = 0;
int k = 0;
for (int i = 0; i < counter; ++i)
{
for (int j = 0; j < counter2; ++j)
{
if (j == i)
{
continue;
}
if (++k == counter3)
{
break;
}
}
if (i == counter2)
{
continue;
}
++accumulate;
}
data[index] = accumulate;
}
/// <summary>
/// Initializes a new instance of the <see cref="MultidirectionalDynamicArray1D{TValue}"/> class.
/// </summary>
/// <param name="dimensions">The dimensions of the array.</param>
public MultidirectionalDynamicArray1D(Index1D dimensions)
{
Contracts.Requires.That(dimensions.IsAllPositive());
this.array = new T[dimensions.X];
this.xOrigin = dimensions.X / 2;
}
/// <summary>
/// This "kernel" function will compile to IL code which ILGPU will ingest and convert to GPU compute shader code.
///
/// </summary>
/// <param name="index"></param>
/// <param name="displayParams"></param> displayPort int[] {width, height }
/// <param name="viewAreaParams"></param> displayView float[] {h_a, h_b, v_a, v_b}
/// <param name="maxIterations"></param>
/// <param name="output"></param>
public static void MandelbrotKernel(
Index1D index,
ArrayView1D <int, Stride1D.Dense> displayParams, ArrayView1D <float, Stride1D.Dense> viewAreaParams, int maxIterations,
ArrayView <int> output)
{
if (index >= output.Length)
{
return;
}
int img_x = index % displayParams[0];
int img_y = index / displayParams[0];
float x0 = viewAreaParams[0] + img_x * (viewAreaParams[1] - viewAreaParams[0]) / displayParams[0];
float y0 = viewAreaParams[2] + img_y * (viewAreaParams[3] - viewAreaParams[2]) / displayParams[1];
float x = 0.0f;
float y = 0.0f;
int iteration = 0;
while ((x * x + y * y < 2 * 2) && (iteration < maxIterations))
{
float xtemp = x * x - y * y + x0;
y = 2 * x * y + y0;
x = xtemp;
iteration += 1;
}
output[index] = iteration;
}
public static void Set <T>(this T[] array, Index1D index, T value)
{
Contracts.Requires.That(array != null);
Contracts.Requires.That(array.IsIndexValid(index));
array[index.X] = value;
}
/// <summary>
/// Handles the array resizing if necessary when an index is accessed.
/// </summary>
/// <param name="index">The index being accessed.</param>
/// <returns>The adjusted index to be used for accessing the backing array.</returns>
private Index1D HandleArrayResizing(Index1D index)
{
int x = index.X;
if (this.IsIndexInCurrentBounds(index))
{
x += this.xOrigin;
return(new Index1D(x));
}
// determine size of new array
int xOffset;
int xNewSize = DynamicArrayUtilities.HandleAxis(
this.GetCurrentLength(Axis1D.X), ref x, ref this.xOrigin, out xOffset);
// copy old array into new array
T[] newArray = new T[xNewSize];
foreach (var pair in this.array.GetIndexValuePairs())
{
newArray[pair.Key.X + xOffset] = pair.Value;
}
this.array = newArray;
return(new Index1D(x));
}
public static T Get <T>(this T[] array, Index1D index)
{
Contracts.Requires.That(array != null);
Contracts.Requires.That(array.IsIndexValid(index));
return(array[index.X]);
}
public TileInfo(ArrayView <T> input, Index1D numIterationsPerGroup)
{
TileSize = Group.DimX * numIterationsPerGroup;
StartIndex = Grid.IdxX * TileSize + Group.IdxX;
EndIndex = (Grid.IdxX + Index1D.One) * TileSize;
MaxLength = XMath.Min(input.IntLength, EndIndex);
}
internal static void EmitRefUsingOutputParamsKernel(
Index1D index,
ArrayView1D <long, Stride1D.Dense> buffer)
{
if (CudaAsm.IsSupported)
{
Input <int> input = index.X;
Input <int> length = buffer.IntLength;
Output <long> result = default;
CudaAsm.EmitRef(
"{\n\t" +
" .reg .s64 t0;\n\t" +
" .reg .s64 t2;\n\t" +
" cvt.s64.s32 t0, %0;\n\t" +
" cvt.s64.s32 t2, %2;\n\t" +
" sub.s64 %1, t2, t0;\n\t" +
"}",
ref input,
ref result,
ref length);
buffer[index] = result.Value;
}
else
{
buffer[index] = index;
}
}
internal static void OffsetOfKernel <T>(
Index1D _,
ArrayView1D <int, Stride1D.Dense> data)
where T : unmanaged
{
data[0] = Interop.OffsetOf <T>("Val2");
}
internal static void WarpShuffleUpKernel(
Index1D index,
ArrayView1D <int, Stride1D.Dense> data,
int shiftAmount)
{
data[index] = Warp.ShuffleUp(Warp.LaneIdx, shiftAmount);
}
internal static void SizeOfKernel <T>(
Index1D _,
ArrayView1D <int, Stride1D.Dense> data)
where T : unmanaged
{
data[0] = Interop.SizeOf <T>();
}
internal static void PrintFKernel(
Index1D index,
ArrayView1D <int, Stride1D.Dense> data,
byte b,
short s,
int i,
long l,
float f,
double d)
{
bool t = b > 0;
Interop.Write(
"Current Thread:\t{0}, {1}+{2} (3) - {3} [{2}]",
(int)index,
b,
s,
i);
Interop.WriteLine(
" => %%:\t{0}, {1}\t{2}/{3}",
l,
f,
d,
t);
}
internal static void BreakKernel(
Index1D index,
ArrayView1D <int, Stride1D.Dense> data,
int counter,
int counter2,
int counter3)
{
int accumulate = 1;
for (int i = 0; i < counter; ++i)
{
if (i == counter2)
{
break;
}
++accumulate;
if (i == counter3)
{
break;
}
++accumulate;
}
data[index] = accumulate;
}
internal static void IfNestedSideEffectsKernel(
Index1D index,
ArrayView1D <int, Stride1D.Dense> data,
ArrayView1D <int, Stride1D.Dense> data2,
int c,
int d)
{
if (c != 0)
{
if (d == 0)
{
data[index] = 42;
data2[index] = 1;
}
else
{
data[index] = 43;
data2[index] = 2;
}
}
else
{
if (d != 0)
{
data[index] = 23;
data2[index] = 3;
}
else
{
data[index] = 24;
data2[index] = 4;
}
}
}
internal static void NestedBreakContinueConstantKernel(
Index1D index,
ArrayView1D <int, Stride1D.Dense> data)
{
int accumulate = 0;
int k = 0;
for (int i = 0; i < 32; ++i)
{
for (int j = 0; j < 13; ++j)
{
if (j == i)
{
continue;
}
if (++k == 9)
{
break;
}
}
if (i == 13)
{
continue;
}
++accumulate;
}
data[index] = accumulate;
}
internal static void NestedIfAndOrKernel(
Index1D index,
ArrayView1D <int, Stride1D.Dense> data,
int c,
int d)
{
int value = d;
if (c < 23)
{
if ((index.X == 0 || index.X == 1) && index.X <= 2)
{
value = 42;
}
if ((index.X == 3 || index.X == 4 || index.X <= 5) && c < 42)
{
value = 43;
}
}
else if (c == 23 || c < 43 && c > d)
{
value = 24;
}
data[index] = value;
}
internal static void NestedCallOutKernel(
Index1D index,
ArrayView1D <int, Stride1D.Dense> data)
{
GetValue(out int value);
data[index] = value;
}
internal static void NestedIfAndOrKernelSideEffects(
Index1D index,
ArrayView1D <int, Stride1D.Dense> data,
int c,
int d)
{
if (c < 23)
{
if ((index.X == 0 || index.X == 1) && index.X <= 2)
{
data[index] = 42;
}
if ((index.X == 3 || index.X == 4 || index.X <= 5) && c < 42)
{
data[index] = 43;
return;
}
}
else if (c == 23 || c < 43 && c > d)
{
data[index] = 24;
return;
}
data[index] = d;
}
internal static void ArrayViewLeaIndexKernel(
Index1D index,
ArrayView1D <int, Stride1D.Dense> data,
ArrayView1D <int, Stride1D.Dense> source)
{
data[index] = source[index];
}
/// <summary>
/// Generate random numbers within a kernel.
/// </summary>
public static void MyRandomKernel(
Index1D index,
RNGView <XorShift64Star> rng,
ArrayView1D <long, Stride1D.Dense> view)
{
view[index] = rng.NextLong();
}
internal static void BasicNestedLoopJumpKernel(
Index1D index,
ArrayView1D <int, Stride1D.Dense> data,
ArrayView1D <int, Stride1D.Dense> source,
int c)
{
int k = 0;
entry:
for (int i = 0; i < source.Length; ++i)
{
if (source[i] == 23)
{
if (k < c)
{
goto exit;
}
goto nested;
}
}
data[index] = 42;
return;
nested:
k = 43;
exit:
if (k++ < 1)
{
goto entry;
}
data[index] = 23 + k;
}
internal static void DoWhileIncrementKernel(
Index1D index,
ArrayView1D <int, Stride1D.Dense> data,
ArrayView1D <int, Stride1D.Dense> data2,
ArrayView1D <int, Stride1D.Dense> data3)
{
int value = 3;
int counter = 0;
do
{
++value;
}while (counter++ < 3);
data[index] = value;
int value2 = 1;
do
{
value2 *= 3;
}while (counter++ == 5);
data2[index] = value2;
int value3 = 13;
do
{
--value3;
}while (counter++ != 8);
data3[index] = value3;
}
internal static void DebugAssertMessageKernel(
Index1D index,
ArrayView1D <int, Stride1D.Dense> data)
{
Debug.Assert(data[index] >= 0, "Invalid kernel argument");
Trace.Assert(data[index] >= 0, "Invalid kernel argument");
}
internal static void ArrayViewGetVariableViewKernel(
Index1D index,
ArrayView1D <int, Stride1D.Dense> data,
ArrayView1D <int, Stride1D.Dense> source)
{
data[index] = source.VariableView(index).Value;
}
