metal - 从内核检索结果

Question

我在玩内核函数时遇到了一些问题。

我想要做的只是简单地向函数发送一个数组，然后在数组中的 waitUntilCompleted 之后获取结果。

以下是一个数组，它将在循环中的 malloc 之后填充从 0 到 123455 的数字:

float *myVector = malloc(123456 * sizeof(float));

这是将发送到内核的数组以及 myVector :

float *resultData =  malloc(123456 * sizeof(float));
id <MTLBuffer> inBuffer = [device newBufferWithBytes:&myVector[0] length:sizeof(myVector) options:MTLResourceOptionCPUCacheModeDefault];
id <MTLBuffer> buffer = [device newBufferWithBytes:&resultData[0] length:sizeof(resultData) options:MTLResourceOptionCPUCacheModeDefault];

使用计算命令编码器，它们都分别设置为索引 0 、 1 和偏移量 0 。

以下设置线程组和组内线程的大小:

MTLSize threadGroupCounts = MTLSizeMake([device maxThreadsPerThreadgroup].width, 1, 1);
MTLSize threadGroups = MTLSizeMake((123456) / threadGroupCounts.width, 1, 1);

[commandEncoder dispatchThreadgroups:threadGroups threadsPerThreadgroup:threadGroupCounts];

[commandEncoder endEncoding];
[commandBuffer commit];
[commandBuffer waitUntilCompleted];

我收到两次以下错误:

Execution of the command buffer was aborted due to an error during execution. Caused GPU Hang Error (IOAF code 3)

在花了很多时间之后，我得出的结论是错误是由以下几行引起的:

MTLSize threadGroupCounts = MTLSizeMake([device maxThreadsPerThreadgroup].width, 1, 1);
MTLSize threadGroups = MTLSizeMake((123456) / [device maxThreadsPerThreadgroup].width, 1, 1);

例如，如果我将 (123456) / [device maxThreadsPerThreadgroup].width 设置为 32 ，则不会发生错误，但除数组中的前 2 个值外，结果将全部为零。

以下是我尝试在处理后获得结果的方法:

NSData *data = [NSData dataWithBytesNoCopy:buffer.contents length:sizeof(myVector) freeWhenDone:NO];
float *finalArray = malloc(sizeof(float) * 123456);
[data getBytes:&finalArray[0] length:sizeof(finalArray)];

这是函数:

kernel void test(const device float *inVector [[buffer (0)]],
                 device float *outVector [[buffer (1)]],
                 uint id [[thread_position_in_grid]])
{
    outVector[id] = -inVector[id]; 
}

我想我在设置线程尺寸时遇到了麻烦。
作为测试，我试图实现的是设置每个线程组允许的最大线程数，将数组的大小除以这个数字并将其发送给处理。有人可以告诉我如何设置线程组大小，将数组发送到函数并最终正确正确地检索数组中的结果吗？

谢谢。

Answer 1

您计算 MTLBuffer 的大小的方式有误。因为 myVector 是一个指针，所以 sizeof(myVector) 可能是 8，而不是 493824。这反过来又会导致您没有为数据分配足够的空间，并且读取超出内核函数中缓冲区的边界。创建缓冲区时尝试使用与使用 malloc 分配浮点数组时相同的大小，看看是否有帮助。

您需要使用 getBytes:length: 对从输出缓冲区中检索的字节数进行相应的更改。

我认为您计算线程组大小和计数的方式是合理的，但您应该注意整数截断。如果要处理的元素总数不能被线程组大小整除，则计算线程组计数的方式将向下取整，从而导致您跳过某些元素。

避免这种情况的一种方法是将您调度的线程组的数量取整，并明确检查缓冲区长度以防止越界访问。所以你可以像这样计算你的线程组数量和大小:

const int elementCount = 123456;
MTLSize threadgroupSize = MTLSizeMake([device maxThreadsPerThreadgroup].width, 1, 1);
MTLSize threadgroups = MTLSizeMake(ceil(elementCount / (float)threadgroupSize.width), 1, 1);

...传入缓冲区大小，如下所示:

[computeCommandEncoder setBytes:&elementCount length:sizeof(elementCount) atIndex:2];

...并检查这样的界限:

kernel void test(const device float *inVector [[buffer (0)]],
                 device float *outVector [[buffer (1)]],
                 constant int &elementCount [[buffer (2)]],
                 uint id [[thread_position_in_grid]])
{
    if (id < elementCount) {
        outVector[id] = -inVector[id];
    }
}