c++ - 我的 iOS Metal 计算内核是否存在编译器错误，或者我遗漏了什么？

Question

我需要一个 upper_bound 的实现如我的 Metal 计算内核的 STL 中所述。 Metal 标准库中没有任何东西，我基本上是从 <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]永远不会被写入——因为我用一个魔数(Magic Number)预加载了缓冲区，看看它是否被覆盖了。它没有。事实上在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 有同样的问题——然而，一个天真的线性版本是有效的(某种程度上)。顺便说一句，我测试了我从直接 C 代码复制的 STL 版本(明显删除了 device)，它们在着色器领域之外工作正常。请告诉我我做错了什么，这不是 Metal 编译器错误。

现在关于上面的“排序”:线性搜索版本在 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);
}

这完全奏效了。也就是说，它不仅解决了线性搜索的“排序”问题，而且类似重写的二分搜索也起作用了。我不想相信这一点。 Metal 着色器语言应该是 C++ 的子集，但标准指针语义不起作用？我真的不能比较或减去指针吗？

无论如何，我不记得看到任何文档说不能有指针别名或者什么声明咒语可以帮助我。还有其他帮助吗？

[更新]

郑重声明，正如 Apple 开发论坛上的“slime”所指出的: 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() 所做的是创建一个本地临时对象，该临时对象指向与其参数之一相同的缓冲区。也许文档应该说出它的意思，比如:“不允许在任何函数中使用指向设备和常量内存的指针别名，包括右值。”我实际上不知道这是否太强了。