Как суммировать FloatResidentArray и получить значение на устройство или хост

Question

Я использую Hybridizer для суммирования FloatResidentArray, и я не могу вернуть рассчитанную сумму на устройство (или хост) из-за необходимости в операторе ref в заключительном операторе AtomicExpr.apply. Рассмотрим следующий код, который основан на примере GenericReduce, предоставленном Altimesh. Код принимает резидентный массив устройства a с плавающей запятой длины N и вычисляет сумму - это значение помещается в 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);
     }
  }

Приведенный выше код не компилируется, поскольку вы не можете передать float [] и FloatResidentArray в одном и том же списке параметров.

Если total определен как сам FloatResidentArray, то компилятор не позволит использовать ключевое слово ref в последней строке кода.

Если я просто передам число с плавающей запятой, возвращаемая переменная не будет обновлена ​​с суммой.

Если я передам ref float - тогда программа выдаст ошибку времени выполнения в тот момент, когда HybRunner оборачивает приведенный выше код для создания динамического - сообщение об ошибке будет

Типы значений по ссылке не поддерживаются

Как мне вернуть сумму? - либо на устройство, либо на хост-память - оба приемлемы.

Answer 1

Что ж, вам нужно понять, как работает сортировка

Объект и массивы (даже резидентный массив) - все хосты при создании в .Net.Затем мы выполняем их маршализацию (закрепление памяти хоста, выделение памяти устройства и копирование хоста на устройство) непосредственно перед выполнением ядра.

• Для float [] это будет сделано автоматически
• Для IntPtr мы ничего не делаем, и пользователь должен убедиться, что IntPtr является действительным указателем устройства, содержащим данные
• Для резидентного массива мы ничего не делаем, и пользователь должен вручную вызвать RefreshDevice () и RefreshHost, когда он хочет получить данные взад и вперед.

Поддерживается смешивание ResidentArray и float [], как показано на снимке экрана сгенерированного dll:

Что не поддерживается: смешиваниеуправляемые типы и IntPtr.

Вот полная версия вашего кода, работающая и возвращающая правильный результат:

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]);

        }
    }
}