Metal 计算值与 CPU 值不同

Question

我正在尝试实现 3DPoints 向量的长度，当我将 GPU 检索到的值与 CPU 进行比较时，它们并不完全相同，通常存在大量差异。 我最初使用了 packed_float3，它存在更多差异，所以我开始使用 float3 并进行了一些改进，但仍有一些差异我想修复。

这些值差别不大，平均而言它们相差 -0.00000000048358334004，但是当我运行诸如对两个数组求和和相减之类的操作时，差异不会发生，我希望它会发生相同的情况。

这是代码的一部分

main.m

- (void) lenght_function:(NSArray*) array {
    _buffer[0] = [_mDevice newBufferWithLength:_sp_size_alloc options:MTLResourceStorageModeShared];
    _buffer[1] = [_mDevice newBufferWithLength:_sp_size_alloc options:MTLResourceStorageModeShared];
    float3 *datapt = [_buffer[0] contents];

    for (unsigned long index = 0 ; index< _sp_lenght ; index++) {
        datapt[index].x = (float)[array[index] getX];
        datapt[index].y = (float)[array[index] getY];
        datapt[index].z = (float)[array[index] getZ];


    }
    commandBuffer = [_mCommandQueue commandBuffer];
    assert(commandBuffer != nil);
    
    id<MTLComputeCommandEncoder> computeEncoder = [commandBuffer computeCommandEncoder];
    assert(computeEncoder != nil);

    [computeEncoder setComputePipelineState:_mLenghtFunctionPSO];
    
    [computeEncoder setBuffer:_buffer[0] offset:0 atIndex:0];
    [computeEncoder setBuffer:_buffer[1] offset:0 atIndex:1];
    //[array1 makeData];
    
    MTLSize gridSize = MTLSizeMake(_sp_lenght, 1, 1);
   
    NSUInteger threadGroupSize = _mLenghtFunctionPSO.maxTotalThreadsPerThreadgroup;
    if(threadGroupSize > _sp_lenght){
        threadGroupSize = _sp_lenght;
    }
    
    MTLSize threadgroupsize = MTLSizeMake(threadGroupSize, 1, 1);
    
    [computeEncoder dispatchThreads:gridSize threadsPerThreadgroup:threadgroupsize];
    [computeEncoder endEncoding];
    [commandBuffer commit];
    [commandBuffer waitUntilCompleted];
    float3 *arr1 = _buffer[0].contents;
    float* result = _buffer[1].contents;
    unsigned long counter = 0;
    for (unsigned long index = 0; index < _sp_lenght; index++)
    {
        if (result[index] != sqrtf(arr1[index].x*arr1[index].x + arr1[index].y*arr1[index].y + arr1[index].z*arr1[index].z)){
            counter++;;
        }
    }
    NSLog(@"ERROR counter %lu\n",counter);

}

kernel.metal

kernel void lenght(const device float3 *arr1,
                         device float *result,
                         uint index[[thread_position_in_grid]]){
    
    result[index] = precise::sqrt(precise::pow(arr1[index].x,2) + precise::pow(arr1[index].y,2) + precise::pow(arr1[index].z,2));

}

Answer 1

32 位精度只有大约 7 个小数位，而您显示的差异大约是 9-10 个小数位。因此，您所展示的内容实际上比人们期望的 32 位浮点精度要好一些。听起来您想要 64 位双精度，但这不是内置的 Metal 数据类型。

如果您将这些值乘以 100 或 1000 以将小数位向上移动，然后在您的值相加后除以该数字，这可能会有所帮助。

另一种可能性是首先标准化你的值，所以它们都在 0 到 1 的范围内。然后你甚至可以使用半精度。