How to total a FloatResidentArray and retrieve the value to the device or host

user_11658290 0
Viewed 54

I am using Hybridizer to total a FloatResidentArray and I am not able to return the calculated total to the device (or host) because of the need for a ref statement in the final AtomicExpr.apply statement. Consider the following code which is based on the GenericReduce example provided by Altimesh. The code takes a device resident array a, of float of length N and calculates the total – this value is placed in total[0].

[Kernel]
 public static void Total(FloatResidentArray a, int N, float[] total)
 {
    var cache = new SharedMemoryAllocator<float>().allocate(blockDim.x);

    int tid = threadIdx.x + blockDim.x * blockIdx.x;
    int cacheIndex = threadIdx.x;
    float sum = 0f;           
    while (tid < N)
    {
       sum = sum + a[tid];               
       tid += blockDim.x * gridDim.x;
     }
     cache[cacheIndex] = sum;          
     CUDAIntrinsics.__syncthreads();
     int i = blockDim.x / 2;
     while (i != 0)
     {
        if (cacheIndex < i)
        {
            cache[cacheIndex] = cache[cacheIndex] + cache[cacheIndex + i];
        }
        CUDAIntrinsics.__syncthreads();
        i >>= 1;
     }

     if (cacheIndex == 0)
     {
          AtomicExpr.apply(ref total[0], cache[0], (x, y) => x + y);
     }
  }

The above code does not compile because you cannot pass a float[] and a FloatResidentArray in the same parameter list.

If total is defined as a FloatResidentArray itself, then the compiler will not allow the ref keyword to be used in the final line of code.

If I simply pass a float, then the returned variable is not updated with the total.

If I pass a ref float - then the program throws a runtime error at the point where the HybRunner wraps the above code to create the dynamic – the error message is

Value types by reference are not supported

How do I return the total? –either to Device or Host memory – both are acceptable.

hybridizer c#
1 Answers

Well, you need to understand how marshalling works

Object and arrays (even resident array) are all hosts when created in .Net. Then we marshal them (pin host memory, allocate device memory and copy host to device) right before kernel execution.

  • For a float[], that will be done automatically
  • For an IntPtr, we do nothing and the user has to ensure the IntPtr is a valid device pointer containing the data
  • For a resident array, we do nothing and the user has to manually call RefreshDevice() and RefreshHost when she wants to get the data back and forth.

Mixing ResidentArray and float[] is supported, as show in this screenshot of the generated dll : enter image description here

What is not supported is : mixing managed types and IntPtr.

Here is a complete version of your code working, and returning the correct result: 

using Hybridizer.Runtime.CUDAImports;
using System;
using System.Runtime.InteropServices;

namespace SimpleMetadataDecorator
{
    class Program
    {
        [EntryPoint]
        public static void Total(FloatResidentArray a, int N, float[] total)
        {
            var cache = new SharedMemoryAllocator<float>().allocate(blockDim.x);

            int tid = threadIdx.x + blockDim.x * blockIdx.x;
            int cacheIndex = threadIdx.x;
            float sum = 0f;
            while (tid < N)
            {
                sum = sum + a[tid];
                tid += blockDim.x * gridDim.x;
            }
            cache[cacheIndex] = sum;
            CUDAIntrinsics.__syncthreads();
            int i = blockDim.x / 2;
            while (i != 0)
            {
                if (cacheIndex < i)
                {
                    cache[cacheIndex] = cache[cacheIndex] + cache[cacheIndex + i];
                }
                CUDAIntrinsics.__syncthreads();
                i >>= 1;
            }

            if (cacheIndex == 0)
            {
                AtomicExpr.apply(ref total[0], cache[0], (x, y) => x + y);
            }
        }
        static void Main(string[] args)
        {

            const int N = 1024 * 1024 * 32;
            FloatResidentArray arr = new FloatResidentArray(N);
            float[] res = new float[1];
            for (int i = 0; i < N; ++i)
            {
                arr[i] = 1.0F;
            }

            arr.RefreshDevice();
            var runner = HybRunner.Cuda();
            cudaDeviceProp prop;
            cuda.GetDeviceProperties(out prop, 0);
            runner.SetDistrib(16 * prop.multiProcessorCount, 1, 128, 1, 1, 128 * sizeof(float));
            var wrapped = runner.Wrap(new Program());
            runner.saveAssembly();
            cuda.ERROR_CHECK((cudaError_t)(int)wrapped.Total(arr, N, res));
            cuda.ERROR_CHECK(cuda.DeviceSynchronize());
            Console.WriteLine(res[0]);

        }
    }
}
user_16693700
Related