我的iOS金属计算内核是否存在编译器错误，还是缺少了什么？

Question

我需要一个upper_bound的实现，就像我的金属计算内核在STL中描述的那样。在金属标准库中没有任何内容，我基本上将它从<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;
}

这件工作(算是)。我对来自lower_bound的二进制搜索<algorithm>也有同样的问题--但是一个简单的线性版本可以工作(差不多)。顺便说一句，我测试了我从直接C代码复制的STL版本(很明显地删除了device )，它们在着色器之外运行得很好。请告诉我，我做错了什么，这不是金属编译器的错误。

现在，关于上面的“排序”：线性搜索版本工作在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;

Answer 1

嗯，我找到了解决办法/解决办法。我猜想这是一个指针混叠问题，因为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++的一个子集，但是标准指针语义不起作用吗？我真的不能比较或减去指针吗？

不管怎么说，我不记得有任何文档说不能有指针混叠，或者什么声明咒语可以帮助我在这里。还有什么帮助吗？

更新

为了记录在案，苹果开发论坛上的“黏液”指出：ref/doc/uid/TP40014364-CH4-SW3

不能别名指定为图形或内核函数的参数值的缓冲区(设备和常量)--也就是说，作为参数值传递的缓冲区不能与传递给同一图形或内核函数的单独参数的另一个缓冲区重叠。

但是也值得注意的是，upper_bound()不是内核函数，而upper_bound_test()没有传递别名参数。upper_bound_test()所做的是创建一个本地临时的，它指向与其参数之一相同的缓冲区。也许文档应该说它的意思，比如：“不允许指针混叠到设备和任何函数中的常量内存，包括rvalue。”我不知道这是不是太强了。