如何汇总 FloatResidentArray 并将值检索到设备或主机

Question

我正在使用 Hybridizer 计算 FloatResidentArray 的总和，但由于在最终的 AtomicExpr.apply 语句中需要一个 ref 语句，所以我无法将计算出的总和返回给设备（或主机）。考虑以下基于 Altimesh 提供的 GenericReduce 示例的代码。该代码采用长度为N的浮点数的设备驻留数组a并计算总数——该值放在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]);

        }

    }

}