- android - 多次调用 OnPrimaryClipChangedListener
- android - 无法更新 RecyclerView 中的 TextView 字段
- android.database.CursorIndexOutOfBoundsException : Index 0 requested, 光标大小为 0
- android - 使用 AppCompat 时,我们是否需要明确指定其 UI 组件(Spinner、EditText)颜色
我想使用 GPU 内核对体积执行双阈值。我将每个切片的卷作为只读 image2d_t 发送。我的输出体积是一个二进制体积,其中每一位 指定它的相关体素是启用还是禁用。我的内核检查当前像素值是否在下限/上限阈值范围内,并在二进制体积中启用其对应位。
出于调试目的,我暂时对实际检查进行了评论。我只是使用传递的切片 nr 来确定二进制卷位是否应该打开或关闭。前 14 个切片设置为“打开”,其余设置为“关闭”。我还在 CPU 端验证了这段代码,我粘贴在这篇文章底部的代码。代码显示了两条路径,现在对 CPU 进行了注释。
CPU 代码按预期工作,在应用二进制掩码渲染体积后返回以下图像:
使用我的 GPU 内核运行完全相同的逻辑会返回错误的结果(第一个 3D,第二个切片 View ):
这里出了什么问题?我读到 OpenCL 不支持位字段,但据我从 OpenCL 规范中了解到,它确实支持按位运算符。支持我的位逻辑,它从 32 位字中选择正确的位并翻转它,对吗?或者我的简单标志被认为是一个位域。它所做的是从左侧(不是右侧,因此是减法)选择体素 %32 位。
另一件事可能是传递给我内核的 uint 指针与我期望的不同。我假设这是对指针的有效使用并将数据传递到我的内核。应用于内核中“uint* 字”部分的逻辑是由于每行填充字和每片填充行。 CPU 变体确认指针计算逻辑是有效的。
下面;代码
uint wordsPerRow = (uint)BinaryVolumeWordsPerRow(volume.Geometry.NumberOfVoxels);
uint wordsPerPlane = (uint)BinaryVolumeWordsPerPlane(volume.Geometry.NumberOfVoxels);
int[] dims = new int[3];
dims[0] = volume.Geometry.NumberOfVoxels.X;
dims[1] = volume.Geometry.NumberOfVoxels.Y;
dims[2] = volume.Geometry.NumberOfVoxels.Z;
uint[] arrC = dstVolume.BinaryData.ObtainArray() as uint[];
unsafe {
fixed(int* dimPtr = dims) {
fixed(uint *arrcPtr = arrC) {
// pick Cloo Platform
ComputePlatform platform = ComputePlatform.Platforms[0];
// create context with all gpu devices
ComputeContext context = new ComputeContext(ComputeDeviceTypes.Gpu,
new ComputeContextPropertyList(platform), null, IntPtr.Zero);
// load opencl source
StreamReader streamReader = new StreamReader(@"C:\views\pii-sw113v1\PMX\ADE\Philips\PmsMip\Private\Viewing\Base\BinaryVolumes\kernels\kernel.cl");
string clSource = streamReader.ReadToEnd();
streamReader.Close();
// create program with opencl source
ComputeProgram program = new ComputeProgram(context, clSource);
// compile opencl source
program.Build(null, null, null, IntPtr.Zero);
// Create the event wait list. An event list is not really needed for this example but it is important to see how it works.
// Note that events (like everything else) consume OpenCL resources and creating a lot of them may slow down execution.
// For this reason their use should be avoided if possible.
ComputeEventList eventList = new ComputeEventList();
// Create the command queue. This is used to control kernel execution and manage read/write/copy operations.
ComputeCommandQueue commands = new ComputeCommandQueue(context, context.Devices[0], ComputeCommandQueueFlags.None);
// Create the kernel function and set its arguments.
ComputeKernel kernel = program.CreateKernel("LowerThreshold");
int slicenr = 0;
foreach (IntPtr ptr in pinnedSlices) {
/*// CPU VARIANT FOR TESTING PURPOSES
for (int y = 0; y < dims[1]; y++) {
for (int x = 0; x < dims[0]; x++) {
long pixelOffset = x + y * dims[0];
ushort* ushortPtr = (ushort*)ptr;
ushort pixel = *(ushortPtr + pixelOffset);
int BinaryWordShift = 5;
int BinaryWordBits = 32;
if (
(0 <= x) &&
(0 <= y) &&
(0 <= slicenr) &&
(x < dims[0]) &&
(y < dims[1]) &&
(slicenr < dims[2])
) {
uint* word =
arrcPtr + 1 + (slicenr * wordsPerPlane) +
(y * wordsPerRow) +
(x >> BinaryWordShift);
uint mask = (uint)(0x1 << ((BinaryWordBits - 1) - (byte)(x & 0x1f)));
//if (pixel > lowerThreshold && pixel < upperThreshold) {
if (slicenr < 15) {
*word |= mask;
} else {
*word &= ~mask;
}
}
}
}*/
ComputeBuffer<int> dimsBuffer = new ComputeBuffer<int>(
context,
ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.CopyHostPointer,
3,
new IntPtr(dimPtr));
ComputeImageFormat format = new ComputeImageFormat(ComputeImageChannelOrder.Intensity, ComputeImageChannelType.UnsignedInt16);
ComputeImage2D image2D = new ComputeImage2D(
context,
ComputeMemoryFlags.ReadOnly,
format,
volume.Geometry.NumberOfVoxels.X,
volume.Geometry.NumberOfVoxels.Y,
0,
ptr
);
// The output buffer doesn't need any data from the host. Only its size is specified (arrC.Length).
ComputeBuffer<uint> c = new ComputeBuffer<uint>(
context, ComputeMemoryFlags.WriteOnly, arrC.Length);
kernel.SetMemoryArgument(0, image2D);
kernel.SetMemoryArgument(1, dimsBuffer);
kernel.SetValueArgument(2, wordsPerRow);
kernel.SetValueArgument(3, wordsPerPlane);
kernel.SetValueArgument(4, slicenr);
kernel.SetValueArgument(5, lowerThreshold);
kernel.SetValueArgument(6, upperThreshold);
kernel.SetMemoryArgument(7, c);
// Execute the kernel "count" times. After this call returns, "eventList" will contain an event associated with this command.
// If eventList == null or typeof(eventList) == ReadOnlyCollection<ComputeEventBase>, a new event will not be created.
commands.Execute(kernel, null, new long[] { dims[0], dims[1] }, null, eventList);
// Read back the results. If the command-queue has out-of-order execution enabled (default is off), ReadFromBuffer
// will not execute until any previous events in eventList (in our case only eventList[0]) are marked as complete
// by OpenCL. By default the command-queue will execute the commands in the same order as they are issued from the host.
// eventList will contain two events after this method returns.
commands.ReadFromBuffer(c, ref arrC, false, eventList);
// A blocking "ReadFromBuffer" (if 3rd argument is true) will wait for itself and any previous commands
// in the command queue or eventList to finish execution. Otherwise an explicit wait for all the opencl commands
// to finish has to be issued before "arrC" can be used.
// This explicit synchronization can be achieved in two ways:
// 1) Wait for the events in the list to finish,
//eventList.Wait();
//}
// 2) Or simply use
commands.Finish();
slicenr++;
}
}
}
}
还有我的内核代码:
const sampler_t smp = CLK_FILTER_NEAREST | CLK_ADDRESS_CLAMP | CLK_NORMALIZED_COORDS_FALSE;
kernel void LowerThreshold(
read_only image2d_t image,
global int* brickSize,
uint wordsPerRow,
uint wordsPerPlane,
int slicenr,
int lower,
int upper,
global write_only uint* c )
{
int4 coord = (int4)(get_global_id(0),get_global_id(1),slicenr,1);
uint4 pixel = read_imageui(image, smp, coord.xy);
uchar BinaryWordShift = 5;
int BinaryWordBits = 32;
if (
(0 <= coord.x) &&
(0 <= coord.y) &&
(0 <= coord.z) &&
(coord.x < brickSize[0]) &&
(coord.y < brickSize[1]) &&
(coord.z < brickSize[2])
) {
global uint* word =
c + 1 + (coord.z * wordsPerPlane) +
(coord.y * wordsPerRow) +
(coord.x >> BinaryWordShift);
uint mask = (uint)(0x1 << ((BinaryWordBits - 1) - (uchar)(coord.x & 0x1f)));
//if (pixel.w > lower && pixel.w < upper) {
if (slicenr < 15) {
*word |= mask;
} else {
*word &= ~mask;
}
}
}
最佳答案
两个问题:
您已将“c”声明为“write_only”,但使用了“|=”和“&=”运算符,它们是读-修改-写
正如其他发帖者所提到的,如果两个工作项正在访问同一个词,则读取-修改-写入之间存在竞争条件,这将导致错误。原子操作比非原子操作慢得多,因此尽管可能,但不推荐。
我建议将输出扩大 8 倍并使用字节而不是位。这将使您的输出成为只写的,并且还会消除争用,从而消除竞争条件。
或者(如果数据紧凑性或格式很重要)每个工作项一次处理 8 个元素,并将复合 8 位输出写为单个字节。这将是只写的,没有争用,并且仍然具有您的数据紧凑性。
关于c# - GPU 内核中的指针和位运算符,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/16786685/
Or 运算符 对两个表达式进行逻辑“或”运算。 result = expression1 Or expression2 参数 result 任意数值变量。 expression1 任意
Not 运算符 对表达式执行逻辑非运算。 result = Not expression 参数 result 任意数值变量。 expression 任意表达式。 说明 下表显示如何
Is 运算符 比较两个对象引用变量。 result = object1 Is object2 参数 result 任意数值变量。 object1 任意对象名。 object2 任意
\ 运算符 两个数相除并返回以整数形式表示的结果。 result = number1\number2 参数 result 任意数值变量。 number1 任意数值表达式。 numbe
And 运算符 对两个表达式进行逻辑“与”运算。 result = expression1 And expression2 参数 result 任意数值变量。 expression1
运算符(+) 计算两个数之和。 result = expression1 + expression2 参数 result 任意数值变量。 expression1 任意表达式。 exp
我对此感到困惑snippet : var n1 = 5-"4"; var n2 = 5+"4"; alert(n1); alert(n2); 我知道 n1 是 1。那是因为减号运算符会将字符串“4”转
我想我会得到 12,而不是 7。 w++,那么w就是4,也就是100,而w++, w 将是 8,1000;所以 w++|z++ 将是 100|1000 = 1100 将是 12。 我怎么了? int
Xor 运算符 对两个表达式进行逻辑“异或”运算。 result = expression1 Xor expression2 参数 result 任意数值变量。 expression1
Mod 运算符 两个数值相除并返回其余数。 result = number1 Mod number2 参数 result 任意数值变量。 number1 任意数值表达式。 numbe
Imp 运算符 对两个表达式进行逻辑蕴涵运算。 result = expression1 Imp expression2 参数 result 任意数值变量。 expression1 任
Eqv 运算符 执行两个表达式的逻辑等价运算。 result = expression1 Eqv expression2 参数 result 任意数值变量。 expression1 任
我有一个运算符重载的简单数学 vector 类。我想为我的运算符(operator)获取一些计时结果。我可以通过计时以下代码轻松计时我的 +=、-=、*= 和/=: Vector sum; for(s
我是用户定义比较运算符的新手。我正在读一本书,其中提到了以下示例: struct P { int x, y; bool operator、运算符<等),我们
在 SQL 的维基百科页面上,有一些关于 SQL 中 bool 逻辑的真值表。 [1] 维基百科页面似乎来源于 SQL:2003 标准。 等号运算符 (=) 的真值表与 SQL:2003 草案中的 I
我遇到了一个奇怪的 C++ 运算符。 http://www.terralib.org/html/v410/classoracle_1_1occi_1_1_number.html#a0f2780081f
我正在阅读关于 SO 和 answers 中的一个问题,它被提到为: If no unambiguous matching deallocation function can be found, pr
我偶然发现了这个解决方案,但我无法理解其中到底发生了什么。谁能解释一下! 据我了解,它试图通过计算一半的单元格然后将其加倍来计算 a*b 网格中的单元格数量。但是我无法理解递归调用。 请不要建议其他解
Go的基本类型 布尔类型bool 长度:1字节 取值:布尔类型的取值只能是true或者false,不能用数字来表示 整型 通用整型 int / uint(有符号 / 无符号,下面也类似) 长度:根据运
在本教程中,您将学习JavaScript中可用的不同运算符,以及在示例的帮助下如何使用它们。 什么是运算符? 在JavaScript中,运算符是一种特殊符号,用于对运算数(值和变量)执行操作。例如,
我是一名优秀的程序员,十分优秀!