Есть ли ошибка компилятора для моего металлического ядра iOS или я что-то упустил?

Мне нужна реализация 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>— все же работает наивная линейная версия. Кстати, я протестировал свои версии, скопированные на STL, с прямого C-кода (с device удалены, очевидно), и они отлично работают за пределами страны шейдеров. Пожалуйста, скажите мне, что я делаю что-то не так, и это не ошибка металлического компилятора.

Теперь об этой «сортировке» выше: версии с линейным поиском работают на 5 с и мини-2 (A7s) (возвращает индекс 3 в примере выше), но на 6+ (А8) это дает правильный ответ + 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;