我正在使用 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 包裝上面的代碼以創建動態的地方拋出一個運行時錯誤 - 錯誤消息是不支持引用的值類型我如何返回總數？– 無論是設備還是主機內存 – 兩者都是可以接受的。
1 回答
白豬掌柜的
TA貢獻1893條經驗 獲得超10個贊
那么，您需要了解編組的工作原理
在 .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]);
}
}
}
- 1 回答
- 0 關注
- 116 瀏覽
添加回答
舉報
0/150
提交
取消