Speed up nested loops and bitwise operations with Alea GPU

Question

Speed up nested loops and bitwise operations with Alea GPU

I'm trying to use Alea to speed up the program I'm working on, but I need help.

I need to do a lot of bit matching and bitwise operations with values stored in two arrays.

For each element of my first array, I must perform bitwise and operation with each element of my second array, and then count the bit set to 1 from the result and.

If the result is greater than / equal to a certain value, I need to exit the inner one and move on to the next element of my first array.

The first array is usually large, with millions of elements, the second is usually less than 200,000 elements.

Trying to do all these operations in parallel, here is my code:

[GpuManaged]
private long[] Check(long[] arr1, long[] arr2, int limit)
{
    Gpu.FreeAllImplicitMemory(true);
    var gpu = Gpu.Default;
    long[] result = new long[arr1.Length];
    gpu.For(0, arr1.Length, i =>
    {
        bool found = false;
        long b = arr1[i];
        for (int i2 = 0; i2 < arr2.Length; i2++)
        {
            if (LibDevice.__nv_popcll(b & arr2[i2]) >= limit)
            {
                found = true;
                break;
            }
        }
        if (!found)
        {
            result[i] = b;
        }
    });
    return result;
}

, , , , .

, , - , GPU.

, NVIDIA - GeForce GT 740M.

2 , . , .

private static int[] CheckWithKernel(Gpu gpu, int[] arr1, int[] arr2, int limit)
{
    var lp = new LaunchParam(16, 256);
    var result = new int[arr1.Length];
    try
    {
        using (var dArr1 = gpu.AllocateDevice(arr1))
        using (var dArr2 = gpu.AllocateDevice(arr2))
        using (var dResult = gpu.AllocateDevice<int>(arr1.Length))
        {
            gpu.Launch(Kernel, lp, arr1.Length, arr2.Length, dArr1.Ptr, dArr2.Ptr, dResult.Ptr, limit);
            Gpu.Copy(dResult, result);
            return result;
        }
    }
    finally
    {
        Gpu.Free(arr1);
        Gpu.Free(arr2);
        Gpu.Free(result);
    }
}

private static void Kernel(int a1, int a2, deviceptr<int> arr1, deviceptr<int> arr2, deviceptr<int> arr3, int limit)
{
    var iinit = blockIdx.x * blockDim.x + threadIdx.x;
    var istep = gridDim.x * blockDim.x;
    for (var i = iinit; i < a1; i += istep)
    {
        bool found = false;
        int b = arr1[i];
        for (var j = 0; j < a2; j++)
        {
            if (LibDevice.__nv_popcll(b & arr2[j]) >= limit)
            {
                found = true;
                break;
            }
        }
        if (!found)
        {
            arr3[i] = b;
        }
    }
}

+4

c# cuda aleagpu

Giovanni A. 28 . '18 10:52

1

TheGeneral · Accepted Answer · 2018-02-28T11:37:03+0000

Update

, GCHandle.Alloc()

.

http://www.aleagpu.com/release/3_0_3/doc/advanced_features_csharp.html

, .

Memory<T> Gpu.AllocateDevice<T>(int length)
Memory<T> Gpu.AllocateDevice<T>(T[] array)

T . .NET. Memory<T> , IDisposable, , Memory<T> . A Memory<T> , , . Memory<T>.Ptr a deviceptr<T>, GPU . . , .

using (var dArg1 = gpu.AllocateDevice(arg1))
using (var dArg2 = gpu.AllocateDevice(arg2))
using (var dOutput = gpu.AllocateDevice<int>(Length/2))
{           
    // pointer arithmetics to access subset of data
    gpu.Launch(Kernel, lp, dOutput.Length, dOutput.Ptr, dArg1.Ptr + Length/2, dArg2.Ptr + Length / 2);

    var result = dOutput.ToArray();

    var expected = arg1.Skip(Length/2).Zip(arg2.Skip(Length/2), (x, y) => x + y);

    Assert.That(result, Is.EqualTo(expected));
}

. , , , GCHandle.Alloc() GCHandleType.Pinned ( unsafe code)

, *
*
, , *
, Parallel Lambda fiddlier

private unsafe long[] Check(long[] arr1, long[] arr2, int limit)
{   
   Gpu.FreeAllImplicitMemory(true);
   var gpu = Gpu.Default;    
   var result = new long[arr1.Length];

   // Create some pinned memory
   var resultHandle = GCHandle.Alloc(result, GCHandleType.Pinned);
   var arr2Handle = GCHandle.Alloc(result, GCHandleType.Pinned);
   var arr1Handle = GCHandle.Alloc(result, GCHandleType.Pinned);

   // get the addresses
   var resultPtr = (int*)resultHandle.AddrOfPinnedObject().ToPointer();
   var arr2Ptr = (int*)arr2Handle.AddrOfPinnedObject().ToPointer();
   var arr1Ptr = (int*)arr2Handle.AddrOfPinnedObject().ToPointer();

   // i hate nasty lambda statements i always find local methods easier to read    
   void Workload(int i)
   {
      var found = false;    
      var b = *(arr1Ptr + i);

      for (var j = 0; j < arr2.Length; j++)
      {
         if (LibDevice.__nv_popcll(b & *(arr2Ptr + j)) >= limit)
         {
            found = true;
            break;
         }
      }

      if (!found)
      {
         *(resultPtr + i) = b;
      }
   }

   try
   {
      gpu.For(0, arr1.Length, i => Workload(i));
   }
   finally 
   {
      // make sure we free
      arr1Handle.Free();
      arr2Handle.Free();
      resultHandle.Free();
   } 
   return result;    
}