예제 #1
0
파일: ArrayViews.cs 프로젝트: rpfeuti/ILGPU
        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;
        }
예제 #2
0
 /// <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;
 }
예제 #3
0
 internal static void DebugAssertKernel(
     Index1D index,
     ArrayView1D <int, Stride1D.Dense> data)
 {
     Debug.Assert(data[index] >= 0);
     Trace.Assert(data[index] >= 0);
 }
예제 #4
0
 internal static void NestedStructureCallOutKernel(
     Index1D index,
     ArrayView1D <int, Stride1D.Dense> data)
 {
     GetStructureValue(out Parent value);
     data[index] = value.Second.Value;
 }
예제 #5
0
        /// <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;
        }
예제 #6
0
        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];
        }
예제 #7
0
파일: BasicLoops.cs 프로젝트: rpfeuti/ILGPU
        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;
        }
예제 #9
0
        /// <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;
        }
예제 #10
0
    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));
        }
예제 #12
0
    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]);
    }
예제 #13
0
 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);
 }
예제 #14
0
        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;
            }
        }
예제 #15
0
 internal static void OffsetOfKernel <T>(
     Index1D _,
     ArrayView1D <int, Stride1D.Dense> data)
     where T : unmanaged
 {
     data[0] = Interop.OffsetOf <T>("Val2");
 }
예제 #16
0
 internal static void WarpShuffleUpKernel(
     Index1D index,
     ArrayView1D <int, Stride1D.Dense> data,
     int shiftAmount)
 {
     data[index] = Warp.ShuffleUp(Warp.LaneIdx, shiftAmount);
 }
예제 #17
0
 internal static void SizeOfKernel <T>(
     Index1D _,
     ArrayView1D <int, Stride1D.Dense> data)
     where T : unmanaged
 {
     data[0] = Interop.SizeOf <T>();
 }
예제 #18
0
        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);
        }
예제 #19
0
파일: BasicLoops.cs 프로젝트: rpfeuti/ILGPU
        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;
        }
예제 #20
0
 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;
         }
     }
 }
예제 #21
0
파일: BasicLoops.cs 프로젝트: rpfeuti/ILGPU
        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;
        }
예제 #22
0
        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;
        }
예제 #23
0
 internal static void NestedCallOutKernel(
     Index1D index,
     ArrayView1D <int, Stride1D.Dense> data)
 {
     GetValue(out int value);
     data[index] = value;
 }
예제 #24
0
 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;
 }
예제 #25
0
 internal static void ArrayViewLeaIndexKernel(
     Index1D index,
     ArrayView1D <int, Stride1D.Dense> data,
     ArrayView1D <int, Stride1D.Dense> source)
 {
     data[index] = source[index];
 }
예제 #26
0
 /// <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();
 }
예제 #27
0
파일: BasicJumps.cs 프로젝트: m4rs-mt/ILGPU
        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;
        }
예제 #28
0
        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;
        }
예제 #29
0
 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");
 }
예제 #30
0
파일: ArrayViews.cs 프로젝트: rpfeuti/ILGPU
 internal static void ArrayViewGetVariableViewKernel(
     Index1D index,
     ArrayView1D <int, Stride1D.Dense> data,
     ArrayView1D <int, Stride1D.Dense> source)
 {
     data[index] = source.VariableView(index).Value;
 }