У меня есть функция ядра в Metal, которой я передаю текстуру, чтобы я мог выполнять некоторые операции с изображением. Я передаю uint2 gid [[thread_position_in_grid]], который дает мне координаты пикселей в виде целых чисел.
Чтобы получить нормализованные координаты устройств, я могу проделать простую математику с gid.x и gid.y вместе с шириной и высотой текстуры. Это лучший способ сделать это? Лучший способ?
Ваш подход хорош. Если вы не хотите запрашивать размеры текстуры внутри функции ядра или создавать буфер только для их передачи, вы можете использовать метод -[MTLComputeCommandEncoder setBytes:length:atIndex:] для привязки размеров текстуры во «временном» буфере, который обрабатывается Metal:
[computeEncoder setBytes:&dimensions length:sizeof(dimensions) atIndex:0]
Думаю, вы правы, и это хороший способ использовать тот же подход, который обычно применяется в GLSL:
float2 texSize = float2(1/outTexture.get_with(),1/outTexture.get_height());
constexpr sampler s(address::clamp_to_edge, filter::linear, coord::normalized);
//
// something to do...
//
float4 color = inTexture.sample(s,float2(gid)*texSize);
//
// something todo with pixel
//
outTexture.write(color,gid);
Метод, указанный в вопросе, работает хорошо. Но для завершения альтернативным способом чтения текстур с использованием ненормализованных (и / или нормализованных координат устройства) было бы использование семплеров.
Создайте сэмплер:
id<MTLSamplerState> GetSamplerState()
{
MTLSamplerDescriptor *desc = [[MTLSamplerDescriptor alloc] autorelease];
desc.minFilter = MTLSamplerMinMagFilterNearest;
desc.magFilter = MTLSamplerMinMagFilterNearest;
desc.mipFilter = MTLSamplerMipFilterNotMipmapped;
desc.maxAnisotropy = 1;
desc.sAddressMode = MTLSamplerAddressModeClampToEdge;
desc.tAddressMode = MTLSamplerAddressModeClampToEdge;
desc.rAddressMode = MTLSamplerAddressModeClampToEdge;
// The key point: specifies that the sampler reads non-normalized coordinates
desc.normalizedCoordinates = NO;
desc.lodMinClamp = 0.0f;
desc.lodMaxClamp = FLT_MAX;
id <MTLSamplerState> sampler_state = nil;
sampler_state = [[device_ newSamplerStateWithDescriptor:desc] autorelease];
// Release the descriptor
desc = nil;
return sampler_state;
}
А затем прикрепите его к кодировщику вычислительной команды:
id <MTLComputeCommandEncoder> compute_encoder = [comand_buffer computeCommandEncoder];
id<MTLSamplerState> ss = GetSamplerState();
// Attach the sampler state to the encoder, say at sampler bind point 0
[compute_encoder setSamplerState:ss atIndex:0];
// And set your texture, say at texture bind point 0
[compute_encoder setTexture:my_texture atIndex:0];
Наконец, используйте его в ядре:
// An example kernel that samples from a texture,
// writes one component of the sample into an output buffer
kernel void compute_main(
texture2d<uint, access::sample> tex_to_sample [[ texture(0) ]],
sampler smp [[ sampler(0) ]],
device uint *out [[buffer(0)]],
uint2 tid [[thread_position_in_grid]])
{
out[tid] = tex_to_sample.sample(smp, tid).x;
}
Использование сэмплера позволяет вам указать параметры для сэмплирования (например, фильтрацию). Вы также можете получить доступ к текстуре по-разному, используя разные семплеры, подключенные к одному и тому же ядру. Sampler также избегает необходимости проходить и проверять границы размеров текстуры.
Обратите внимание, что сэмплер также можно настроить из вычислительного ядра. См. Раздел 2.6 «Сэмплеры» в Спецификации языка затенения металлов
Наконец, одно из основных различий между функцией чтения (с использованием gid, как указано в вопросе) и выборкой с использованием сэмплера заключается в том, что read () принимает целочисленные координаты, тогда как sample () принимает координаты с плавающей запятой. Таким образом, целочисленные координаты, переданные в образец, будут преобразованы в эквивалентные числа с плавающей запятой.
