metal - Metal Compute Kernel 因两个纹理参数而失败

Question

我有金属计算内核，它需要两个纹理作为参数。但是，我遇到了内核无法运行的问题。我已将问题归结为这个简单的内核。

#include <metal_stdlib>
using namespace metal;

kernel void test_texture(texture2d<float, access::sample>  tex1 [[texture(0)]],
                         texture2d<float, access::sample>  tex2 [[texture(1)]],
                         device float                      *buf  [[buffer(0)]],
                         uint                              idx  [[thread_position_in_grid]])
{
    buf[idx] = 100;
}

以及以下主机代码。

#import <Metal/Metal.h>

int main(int argc, const char * argv[]) {
    @autoreleasepool {
        const size_t max_buffer = 128000000;
        const size_t max_texture = 16384;

        id<MTLDevice> device = MTLCreateSystemDefaultDevice();
        id<MTLLibrary> library = [device newDefaultLibrary];
        id<MTLCommandQueue> queue = [device newCommandQueue];

        id<MTLBuffer> buffer = [device newBufferWithLength:sizeof(float)*max_buffer
                                                   options:MTLResourceCPUCacheModeDefaultCache |
                                                           MTLResourceStorageModeManaged];

        MTLTextureDescriptor *textureDescriptor = [[MTLTextureDescriptor alloc] init];
        textureDescriptor.textureType = MTLTextureType2D;
        textureDescriptor.pixelFormat = MTLPixelFormatR32Float;
        textureDescriptor.width = max_texture;
        textureDescriptor.height = max_texture;
        textureDescriptor.depth = 1;
        textureDescriptor.mipmapLevelCount = 1;
        textureDescriptor.sampleCount = 1;
        textureDescriptor.arrayLength = 1;
        textureDescriptor.resourceOptions = MTLResourceStorageModePrivate | MTLResourceCPUCacheModeDefaultCache;
        textureDescriptor.cpuCacheMode = MTLCPUCacheModeDefaultCache;
        textureDescriptor.storageMode = MTLStorageModePrivate;
        textureDescriptor.usage = MTLTextureUsageShaderRead;

        id<MTLTexture> texture1 = [device newTextureWithDescriptor:textureDescriptor];
        id<MTLTexture> texture2 = [device newTextureWithDescriptor:textureDescriptor];

        MTLComputePipelineDescriptor *discriptor = [[MTLComputePipelineDescriptor alloc] init];
        discriptor.computeFunction = [library newFunctionWithName:@"test_texture"];
        discriptor.threadGroupSizeIsMultipleOfThreadExecutionWidth = YES;

        id<MTLComputePipelineState> pipeline = [device newComputePipelineStateWithDescriptor:discriptor
                                                                                     options:MTLPipelineOptionNone
                                                                                  reflection:NULL
                                                                                       error:NULL];

        id<MTLCommandBuffer> command_buffer = queue.commandBuffer;
        id<MTLComputeCommandEncoder> compute_encoder = [command_buffer computeCommandEncoder];
        [compute_encoder setComputePipelineState:pipeline];
        [compute_encoder setTexture:texture1 atIndex:0];
        [compute_encoder setTexture:texture2 atIndex:1];
        [compute_encoder setBuffer:buffer offset:0 atIndex:0];
        [compute_encoder dispatchThreads:MTLSizeMake(max_buffer, 1, 1) threadsPerThreadgroup:MTLSizeMake(1024, 1, 1)];
        [compute_encoder endEncoding];

        id<MTLBlitCommandEncoder> blit_encoder = [command_buffer blitCommandEncoder];
        [blit_encoder synchronizeResource:buffer];
        [blit_encoder endEncoding];

        [command_buffer commit];
        [command_buffer waitUntilCompleted];

        float *result = (float *)buffer.contents;
        NSLog(@"%f",result[0]);
    }
    return 0;
}

如果我注释掉第二个纹理参数，我会在读取结果缓冲区时得到预期值。但是，当我保持第二个纹理参数不变时，似乎内核没有运行并且结果中的值为零。在 MacOS 上的计算内核中可以采样的纹理数量是否有限制？或者是我在两个纹理中使用最大纹理尺寸引起的问题（我的纹理内存用完了）？强文本

标签： metal

Answer 1

在您的情况下，错误很可能是由于纹理占用了您的整个视频内存预算而发生的。16384 x 16384 * sizeof(float) = 每个纹理 1024mb 的内存。因为您使用的是 MTLStorageModePrivate，所以资源仅存储在视频内存中。