我需要针对金属计算内核的STL中所述的upper_bound实现。金属标准库中没有任何内容，因此我基本上将其从<algorithm>复制到了着色器文件中，如下所示：
static device float* upper_bound( device float* first, device float* last, float val)
{
ptrdiff_t count = last - first;
while( count > 0){
device float* it = first;
ptrdiff_t step = count/2;
it += step;
if( !(val < *it)){
first = ++it;
count -= step + 1;
}else count = step;
}
return first;
}
我创建了一个简单的内核来测试它，如下所示：
kernel void upper_bound_test(
device float* input [[buffer(0)]],
device uint* output [[buffer(1)]]
)
{
device float* where = upper_bound( input, input + 5, 3.1);
output[0] = where - input;
}
该测试的哪个具有硬编码的输入大小和搜索值。我还在框架端对5元素输入缓冲区进行了硬编码，如下所示。我希望该内核返回大于3.1的第一个输入的索引
没用实际上，
output[0]从来没有写过-我在缓冲区中预装了一个魔术数字，看它是否被覆盖。没有。实际上，在waitUntilCompleted之后，commandBuffer.error看起来像这样：Error Domain = MTLCommandBufferErrorDomain
Code = 1
NSLocalizedDescription = "IOAcceleratorFamily returned error code 3"
错误代码3是什么意思？我的内核在有机会完成之前被杀死了吗？
此外，我只尝试了
upper_bound的线性搜索版本，如下所示：static device float* upper_bound2( device float* first, device float* last, float val)
{
while( first < last && *first <= val)
++first;
return first;
}
这一项（排序）。我从
<algorithm>进行二进制搜索lower_bound时遇到了同样的问题-但是，一个简单的线性版本有效（排序）。顺便说一句，我测试了我的STL复制版本，使用直接的C代码（明显删除了device），它们在shader-land之外运行良好。请告诉我我做错了，这不是一个金属编译器错误。现在关于上面的“排序”：线性搜索版本在5s和mini-2（A7s）上工作（在上面的示例中返回索引3），但是在6+（A8）上给出正确的答案+ 2 ^ 31。有没有搞错！完全相同的代码。请注意，在框架方面，我使用
uint32_t；在着色器方面，我使用uint；这是同一件事。还要注意，每个指针减法（ptrdiff_t是带符号的8字节的东西）都是小的非负值。为什么6+设置是那个高阶位？当然，为什么我真正的二进制搜索版本不起作用？这是框架方面的东西：
id<MTLFunction> upperBoundTestKernel = [_library newFunctionWithName: @"upper_bound_test"];
id <MTLComputePipelineState> upperBoundTestPipelineState = [_device
newComputePipelineStateWithFunction: upperBoundTestKernel
error: &err];
float sortedNumbers[] = {1., 2., 3., 4., 5.};
id<MTLBuffer> testInputBuffer = [_device
newBufferWithBytes:(const void *)sortedNumbers
length: sizeof(sortedNumbers)
options: MTLResourceCPUCacheModeDefaultCache];
id<MTLBuffer> testOutputBuffer = [_device
newBufferWithLength: sizeof(uint32_t)
options: MTLResourceCPUCacheModeDefaultCache];
*(uint32_t*)testOutputBuffer.contents = 42;//magic number better get clobbered
id<MTLCommandBuffer> commandBuffer = [_commandQueue commandBuffer];
id<MTLComputeCommandEncoder> commandEncoder = [commandBuffer computeCommandEncoder];
[commandEncoder setComputePipelineState: upperBoundTestPipelineState];
[commandEncoder setBuffer: testInputBuffer offset: 0 atIndex: 0];
[commandEncoder setBuffer: testOutputBuffer offset: 0 atIndex: 1];
[commandEncoder
dispatchThreadgroups: MTLSizeMake( 1, 1, 1)
threadsPerThreadgroup: MTLSizeMake( 1, 1, 1)];
[commandEncoder endEncoding];
[commandBuffer commit];
[commandBuffer waitUntilCompleted];
uint32_t answer = *(uint32_t*)testOutputBuffer.contents;
最佳答案
好吧，我找到了一个解决方案/变通办法。我猜这是指针混淆的问题，因为first和last指向同一个缓冲区。所以我将它们更改为单个指针变量的偏移量。这是重写的upper_bound2：
static uint upper_bound2( device float* input, uint first, uint last, float val)
{
while( first < last && input[first] <= val)
++first;
return first;
}
和一个重写的测试内核：
kernel void upper_bound_test(
device float* input [[buffer(0)]],
device uint* output [[buffer(1)]]
)
{
output[0] = upper_bound2( input, 0, 5, 3.1);
}
这完全起作用了。也就是说，它不仅解决了线性搜索的“排序”问题，而且同样重写了二进制搜索也起作用。我不想相信这一点。金属着色器语言应该是C ++的子集，但是标准指针语义不起作用？我真的可以不比较或减去指针吗？
无论如何，我不记得有任何文档说没有指针别名或什么声明声明会在这里对我有所帮助。还有其他帮助吗？
[更新]
作为记录，正如Apple开发人员论坛上的“史莱姆”所指出的：
https://developer.apple.com/library/ios/documentation/Metal/Reference/MetalShadingLanguageGuide/func-var-qual/func-var-qual.html#//apple_ref/doc/uid/TP40014364-CH4-SW3
“指定为图形或内核函数的参数值的缓冲区（设备和常量）不能混叠-也就是说，作为参数值传递的缓冲区不能与传递给相同图形或内核函数的单独参数的另一个缓冲区重叠。”
但是也值得注意的是，upper_bound（）不是内核函数，并且upper_bound_test（）没有传递别名参数。 upper_bound_test（）的作用是创建一个本地临时文件，该临时文件与其参数之一指向同一缓冲区。也许文档应该说出这是什么意思，例如：“在任何函数中，包括rvalues，都不允许对设备和常量内存使用别名的指针。”我实际上不知道这是否太强了。
关于c++ - 我的iOS Metal 计算内核是否存在编译器错误，或者我缺少某些东西？，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/34585500/