将 UInt8 组件类型的纹理传递给 Metal 计算着色器
Posted
技术标签:
【中文标题】将 UInt8 组件类型的纹理传递给 Metal 计算着色器【英文标题】:Passing textures with UInt8 component type to Metal compute shader 【发布时间】:2018-05-24 02:54:17 【问题描述】:我有一个以编程方式生成的图像,我想将此图像作为纹理发送到计算着色器。我生成此图像的方式是将每个 RGBA 分量计算为UInt8
值,并将它们组合成UInt32
并将其存储在图像的缓冲区中。我使用以下代码来做到这一点:
guard let cgContext = CGContext(data: nil,
width: width,
height: height,
bitsPerComponent: 8,
bytesPerRow: 0,
space: CGColorSpaceCreateDeviceRGB(),
bitmapInfo: RGBA32.bitmapInfo) else
print("Unable to create CGContext")
return
guard let buffer = cgContext.data else
print("Unable to create textures")
return
let pixelBuffer = buffer.bindMemory(to: RGBA32.self, capacity: width * height)
let heightFloat = Float(height)
let widthFloat = Float(width)
for i in 0 ..< height
let latitude = Float(i + 1) / heightFloat
for j in 0 ..< width
let longitude = Float(j + 1) / widthFloat
let x = UInt8(((sin(longitude * Float.pi * 2) * cos(latitude * Float.pi) + 1) / 2) * 255)
let y = UInt8(((sin(longitude * Float.pi * 2) * sin(latitude * Float.pi) + 1) / 2) * 255)
let z = UInt8(((cos(latitude * Float.pi) + 1) / 2) * 255)
let offset = width * i + j
pixelBuffer[offset] = RGBA32(red: x, green: y, blue: z, alpha: 255)
let coordinateConversionImage = cgContext.makeImage()
RGBA32
是一个小结构体,用于移动和创建UInt32
值。这张图片很好,因为我可以将其转换为 UIImage
并将其保存到我的照片库中。
当我尝试将此图像作为纹理发送到计算着色器时,就会出现问题。下面是我的着色器代码:
kernel void updateEnvironmentMap(texture2d<uint, access::read> currentFrameTexture [[texture(0)]],
texture2d<uint, access::read> coordinateConversionTexture [[texture(1)]],
texture2d<uint, access::write> environmentMap [[texture(2)]]
uint2 gid [[thread_position_in_grid]])
const uint4 pixel = 255, 127, 63, 255;
environmentMap.write(pixel, gid);
这段代码的问题是我的纹理类型是uint
,它是32位的,我想像在CPU上一样生成32位像素,通过附加4个8位价值观。但是,我似乎无法在 Metal 上执行此操作,因为没有 byte
类型,我可以将它们附加在一起并组成 uint32
。所以,我的问题是,在 Metal 计算着色器上处理 2D 纹理和设置 32 位像素的正确方法是什么?
额外问题:另外,我看到了以texture2d<float, access::read>
作为输入纹理类型的示例着色器代码。我假设它代表一个介于 0.0 和 1.0 之间的值,但与值介于 0 和 255 之间的无符号整数相比,它有什么优势?
编辑:澄清一下,着色器的输出纹理environmentMap
与输入纹理具有完全相同的属性(宽度、高度、像素格式等)。为什么我认为这是反直觉的,因为我们将uint4
设置为一个像素,这意味着它由 4 个 32 位值组成,而每个像素应该是 32 位。使用此当前代码,255, 127, 63, 255
与2550, 127, 63, 255
具有完全相同的结果,这意味着值在写入输出纹理之前以某种方式被限制在 0-255 之间。但这非常违反直觉。
【问题讨论】:
【参考方案1】:这里有一些你似乎不熟悉的魔法,所以我会试着解释一下。
首先,在设计上,Metal 中纹理的存储格式与读取/采样时获得的类型之间存在松散的联系。您可以拥有.bgra8Unorm
格式的纹理,当通过绑定为texture2d<float, access::sample>
的纹理进行采样时,将为您提供float4
,其组件按RGBA 顺序排列。从这些压缩字节到带有 swizzled 组件的浮点向量的转换遵循 Metal Shading Language Specification 中指定的详细记录的转换规则。
同样的情况是,当写入存储为(例如)每个组件 8 位的纹理时,值将被钳制以适应底层存储格式。这进一步受纹理是否为norm
类型的影响:如果格式包含norm
,则这些值被解释为好像它们指定了一个介于0 和1 之间的值。否则,您读取的值不会被规范化。
一个例子:如果纹理是.bgra8Unorm
,并且给定的像素包含字节值[0, 64, 128, 255]
,那么当在请求float
组件的着色器中读取时,您将在采样时得到[0.5, 0.25, 0, 1.0]
。相反,如果格式为.rgba8Uint
,您将得到[0, 64, 128, 255]
。纹理的存储格式对其内容在采样时的解释方式具有普遍影响。
我假设纹理的像素格式类似于.rgba8Unorm
。如果是这种情况,您可以通过这样编写内核来实现您想要的:
kernel void updateEnvironmentMap(texture2d<float, access::read> currentFrameTexture [[texture(0)]],
texture2d<float, access::read> coordinateConversionTexture [[texture(1)]],
texture2d<float, access::write> environmentMap [[texture(2)]]
uint2 gid [[thread_position_in_grid]])
const float4 pixel(255, 127, 63, 255);
environmentMap.write(pixel * (1 / 255.0), gid);
相比之下,如果你的纹理格式为.rgba8Uint
,你可以这样写得到同样的效果:
kernel void updateEnvironmentMap(texture2d<float, access::read> currentFrameTexture [[texture(0)]],
texture2d<float, access::read> coordinateConversionTexture [[texture(1)]],
texture2d<float, access::write> environmentMap [[texture(2)]]
uint2 gid [[thread_position_in_grid]])
const float4 pixel(255, 127, 63, 255);
environmentMap.write(pixel, gid);
我知道这是一个玩具示例,但我希望通过上述信息,您可以弄清楚如何正确存储和采样值以实现您想要的。
【讨论】:
感谢沃伦,这澄清了很多事情。此外,如果有人想了解这方面的更多信息,可以查看 Metal Shading Language Specification 的第 7.7 章纹理寻址和转换规则。 kCVPixelFormatType_OneComponent8 是规范类型吗?以上是关于将 UInt8 组件类型的纹理传递给 Metal 计算着色器的主要内容,如果未能解决你的问题,请参考以下文章
在大纹理上渲染小纹理时,Metal 比 OpenGL 慢得多
Swift Metal 将 bgra8Unorm 纹理保存到 PNG 文件