是否有我的iOS金屬計算內核的編譯器錯誤，或者我錯過了什麼？

Question

我需要在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），並且它們在着色器之外很好地工作。請告訴我我做錯了，這不是一個金屬編譯器錯誤。現在關於上面的「排序」：線性搜索版本在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++的一個子集，但標準指針語義不起作用？我真的不能比較或減去指針嗎？

無論如何，我不記得看到任何文檔說沒有指針別名或什麼聲明咒語會幫助我在這裏。還有什麼幫助？

[UPDATE]

對於記錄，如蘋果的dev的論壇上指出的「粘液」： 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（）做的是創建一個本地臨時指針，它指向與其參數相同的緩衝區。也許文檔應該說明它的含義，例如：「不允許任何函數的指針別名和任何函數的常量內存，包括右值。」我實際上並不知道這是否太強大。