将 UInt8 组件类型的纹理传递给 Metal 计算着色器

Question

我有一个以编程方式生成的图像，我想将此图像作为纹理发送到计算着色器。我生成此图像的方式是将每个 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 之间。但这非常违反直觉。

Answer 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);
}

我知道这是一个玩具示例，但我希望通过上述信息，您可以弄清楚如何正确存储和采样值以实现您想要的。