我正在使用 Hybridizer 計(jì)算 FloatResidentArray 的總和，但由于在最終的 AtomicExpr.apply 語(yǔ)句中需要一個(gè) ref 語(yǔ)句，所以我無(wú)法將計(jì)算出的總和返回給設(shè)備（或主機(jī)）。考慮以下基于 Altimesh 提供的 GenericReduce 示例的代碼。該代碼采用長(zhǎng)度為N的浮點(diǎn)數(shù)的設(shè)備駐留數(shù)組a并計(jì)算總數(shù)——該值放在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); } }上面的代碼無(wú)法編譯，因?yàn)槟荒茉谕粎?shù)列表中傳遞 float[] 和 FloatResidentArray。如果 total 定義為 FloatResidentArray 本身，那么編譯器將不允許在最后一行代碼中使用 ref 關(guān)鍵字。如果我只是傳遞一個(gè)浮點(diǎn)數(shù)，則返回的變量不會(huì)用總數(shù)更新。如果我傳遞一個(gè) ref float - 然后程序在 HybRunner 包裝上面的代碼以創(chuàng)建動(dòng)態(tài)的地方拋出一個(gè)運(yùn)行時(shí)錯(cuò)誤 - 錯(cuò)誤消息是不支持引用的值類型我如何返回總數(shù)？– 無(wú)論是設(shè)備還是主機(jī)內(nèi)存 – 兩者都是可以接受的。
1 回答
白豬掌柜的
TA貢獻(xiàn)1893條經(jīng)驗(yàn) 獲得超10個(gè)贊
那么，您需要了解編組的工作原理
在 .Net 中創(chuàng)建的對(duì)象和數(shù)組（甚至常駐數(shù)組）都是宿主。然后我們?cè)趦?nèi)核執(zhí)行之前編組它們（固定主機(jī)內(nèi)存，分配設(shè)備內(nèi)存并將主機(jī)復(fù)制到設(shè)備）。
對(duì)于 float[]，這將自動(dòng)完成
對(duì)于 IntPtr，我們什么都不做，用戶必須確保 IntPtr 是包含數(shù)據(jù)的有效設(shè)備指針
對(duì)于常駐數(shù)組，我們什么都不做，用戶在想要來(lái)回獲取數(shù)據(jù)時(shí)必須手動(dòng)調(diào)用 RefreshDevice() 和 RefreshHost。
支持混合 ResidentArray 和 float[]，如生成的 dll 的屏幕截圖所示：
不支持的是：混合托管類型和 IntPtr。
這是您的代碼工作的完整版本，并返回正確的結(jié)果：
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 關(guān)注
- 122 瀏覽
添加回答
舉報(bào)
0/150
提交
取消