我有一个MTLTexture包含16位无符号整数(MTLPixelFormatR16Uint).值范围从大约7000到20000,其中0用作'nodata'值,这就是为什么在下面的代码中跳过它.我想找到最小值和最大值,以便我可以在0-255之间重新调整这些值.最后,我将寻找基于数据直方图的最小值和最大值(它有一些异常值),但是现在我只是简单地提取最小值/最大值.
我可以将GPU中的数据读取到CPU并将最小/最大值拉出,但更愿意在GPU上执行此任务.
第一次尝试
命令编码器被调度为每个线程组16x16个线程,线程组的数量基于纹理大小(例如,width = textureWidth/16,height = textureHeight/16).
typedef struct {
atomic_uint min;
atomic_uint max;
} BandMinMax;
kernel void minMax(texture2d<ushort, access::read> band1 [[texture(0)]],
device BandMinMax &out [[buffer(0)]],
uint2 gid [[thread_position_in_grid]])
{
ushort value = band1.read(gid).r;
if (value != 0) {
uint currentMin = atomic_load_explicit(&out.min, memory_order_relaxed);
uint currentMax = atomic_load_explicit(&out.max, memory_order_relaxed);
if (value > currentMax) {
atomic_store_explicit(&out.max, value, memory_order_relaxed);
}
if (value < currentMin) {
atomic_store_explicit(&out.min, value, memory_order_relaxed);
}
}
}
从中我得到最小值和最大值,但对于相同的数据集,min和max通常会返回不同的值.相当确定这是多线程运行时来自单个线程的最小值和最大值.
第二次尝试
在前一次尝试的基础上,这次我将存储每个线程的各个最小/最大值,全部为256(16x16).
kernel void minMax(texture2d<ushort, …我想使用Metal计算着色器来计算一些位置,然后将其输入到金属着色器中.听起来很直接,但我无法将我的MTLBuffer数据输入基于金属的SCNProgram.
计算内核如下,在这个设计的例子中,它接收三个3D向量(在两个缓冲区中).
kernel void doSimple(const device float3 *inVector [[ buffer(0) ]],
device float3 *outVector [[ buffer(1) ]],
uint id [[ thread_position_in_grid ]]) {
float yDisplacement = 0;
. . . .
outVector[id] = float3(
inVector[id].x,
inVector[id].y + yDisplacement,
inVector[id].z);
}
这个内核函数在- renderer:willRenderScene:atTime:我的方法中每帧运行一次SCNSceneRendererDelegate.有两个缓冲区,它们在每帧之后切换.
缓冲区创建如下;
func setupBuffers() {
positions = [vector_float3(0,0,0), vector_float3(1,0,0), vector_float3(2,0,0)]
let bufferSize = sizeof(vector_float3) * positions.count
//copy same data into two different buffers for initialisation
buffer1 = device.newBufferWithBytes(&positions, length: bufferSize, options: .OptionCPUCacheModeDefault)
buffer2 = device.newBufferWithBytes(&positions, …