前言

Metal入門教程總結(jié)

正文

核心思路

首先坐求，我們用直方圖來表示一張圖像：橫坐標代表的是顏色值伟恶，縱坐標代表的是該顏色值在圖像中出現(xiàn)次數(shù)。

如圖歹苦，對于某些圖像，可能出現(xiàn)顏色值集中分布在某個區(qū)間的情況督怜。
直方圖均衡化(Histogram Equalization) 暂氯，指的是對圖像的顏色值進行重新分配，使得顏色值的分布更加均勻亮蛔。

本文用compute shader對圖像的顏色值進行統(tǒng)計痴施，然后計算得出映射關(guān)系，由fragment shader進行顏色映射處理究流。

效果展示

具體步驟

1辣吃、Metal的render管道、compute管道配置芬探；

同前文神得，不再贅述，詳見Metal入門教程總結(jié)偷仿。

2哩簿、CPU進行直方圖均衡化處理；

• 2.1 把UIImage轉(zhuǎn)成Bytes酝静；

• 2.2 顏色統(tǒng)計节榜；

    // CPU進行統(tǒng)計
    Byte *color = (Byte *)spriteData;
    for (int i = 0; i < width * height; ++i) {
        for (int j = 0; j < LY_CHANNEL_NUM; ++j) {
            uint c = color[i * 4 + j];
            ++cpuColorBuffer.channel[j][c];
        }
    }

• 2.3 映射關(guān)系；

    int rgb[3][LY_CHANNEL_SIZE], sum = (int)(width * height);
    int val[3] = {0};
    // 顏色映射
    for (int i = 0; i < 3; ++i) {
        for (int j = 0; j < LY_CHANNEL_SIZE; ++j) {
            val[i] += cpuColorBuffer.channel[i][j];
            rgb[i][j] = val[i] * 1.0 * (LY_CHANNEL_SIZE - 1) / sum;
        }
    }

• 2.4 顏色值修改别智；

    // 值修改
    for (int i = 0; i < width * height; ++i) {
        for (int j = 0; j < LY_CHANNEL_NUM; ++j) {
            uint c = color[i * 4 + j];
            color[i * 4 + j] = rgb[j][c];
        }
    }

最后用處理之后的Bytes生成新圖片宗苍。

3 GPU進行直方圖均衡化處理；

• 3.1 compute shader進行顏色統(tǒng)計；

kernel void
grayKernel(texture2d<float, access::read>  sourceTexture  [[textureLYKernelTextureIndexSource]], // 紋理輸入讳窟，
           device LYColorBuffer &out [[buffer(LYKernelBufferIndexOutput)]], // 輸出的buffer
           uint2                          grid         [[thread_position_in_grid]]) // 格子索引
{
    // 邊界保護
    if(grid.x < sourceTexture.get_width() && grid.y < sourceTexture.get_height())
    {
        float4 color  = sourceTexture.read(grid); // 初始顏色
        int3 rgb = int3(color.rgb * SIZE); // 乘以SIZE让歼，得到[0, 255]的顏色值
        // 顏色統(tǒng)計，每個像素點計一次
        atomic_fetch_add_explicit(&out.channel[0][rgb.r], 1, memory_order_relaxed);
        atomic_fetch_add_explicit(&out.channel[1][rgb.g], 1, memory_order_relaxed);
        atomic_fetch_add_explicit(&out.channel[2][rgb.b], 1, memory_order_relaxed);
    }
}

atomic_fetch_add_explicit是用于在多線程進行數(shù)據(jù)操作丽啡，具體的函數(shù)解釋見這里谋右。

• 3.2 映射關(guān)系處理；
  compute shader回調(diào)后补箍，根據(jù)GPU統(tǒng)計的顏色分布結(jié)果倚评，求出映射關(guān)系；

        LYLocalBuffer *buffer = (LYLocalBuffer *)strongSelf.colorBuffer.contents; // GPU統(tǒng)計的結(jié)果
        LYLocalBuffer *convertBuffer = self.convertBuffer.contents; // 顏色轉(zhuǎn)換的buffer
        int sum = (int)(self.sourceTexture.width * self.sourceTexture.height); // 總的像素點
        int val[3] = {0}; // 累計和
        for (int i = 0; i < 3; ++i) {
            for (int j = 0; j < LY_CHANNEL_SIZE; ++j) {
                val[i] += buffer->channel[i][j]; // 當前[0, j]累計出現(xiàn)的總次數(shù)
                convertBuffer->channel[i][j] = val[i] * 1.0 * (LY_CHANNEL_SIZE - 1) / sum;
                
                // 對比CPU和GPU處理的結(jié)果
                if (buffer->channel[i][j] != strongSelf->cpuColorBuffer.channel[i][j]) {
                    // 如果不相同馏予，則把對應(yīng)的結(jié)果輸出
                    printf("%d, %d, gpuBuffer:%u  cpuBuffer:%u \n", i, j, buffer->channel[i][j], strongSelf->cpuColorBuffer.channel[i][j]);
                }
            }
        }
        memset(buffer, 0, strongSelf.colorBuffer.length);

3.3 根據(jù)映射關(guān)系處理原圖片天梧，并渲染到屏幕上；

fragment float4
samplingShader(RasterizerData input [[stage_in]], // stage_in表示這個數(shù)據(jù)來自光柵化霞丧。（光柵化是頂點處理之后的步驟呢岗，業(yè)務(wù)層無法修改）
               texture2d<float> colorTexture [[ texture(LYFragmentTextureIndexSource) ]], // texture表明是紋理數(shù)據(jù)，LYFragmentTextureIndexSource是索引
               device LYLocalBuffer &convertBuffer [[buffer(LYFragmentBufferIndexConvert)]]) // 轉(zhuǎn)換的buffer
{
    constexpr sampler textureSampler (mag_filter::linear, min_filter::linear); // sampler是采樣器
    float4 colorSample = colorTexture.sample(textureSampler, input.textureCoordinate); // 得到紋理對應(yīng)位置的顏色
    int3 rgb = int3(colorSample.rgb * SIZE); // 記得先乘以SIZE
    colorSample.rgb = float3(convertBuffer.channel[0][rgb.r], convertBuffer.channel[1][rgb.g], convertBuffer.channel[2][rgb.b]) / SIZE; // 返回的值也要經(jīng)過歸一化處理
    return colorSample;
}

遇到的問題

1蛹尝、統(tǒng)計結(jié)果集中在頭部

問題表現(xiàn)：
統(tǒng)計結(jié)果異常后豫，集中在前面兩個值。
如下突那，green通道的顏色集中在r[0]和r[1]上：

// 0~255顏色值的分布
28269 4492 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0

問題分析：
實際上挫酿，gpu里面存著的是0.0~1.0的值；（歸一化）
統(tǒng)計的值全部是在前面愕难，是因為沒有乘以255早龟！
先用CPU實現(xiàn)了直方圖均衡化，在實現(xiàn)shader的時候猫缭，參考CPU的代碼實現(xiàn)葱弟，犯了這個錯誤。

2猜丹、cpu和gpu統(tǒng)計結(jié)果相差較多

問題表現(xiàn)：
如下代碼芝加，buffer是gpu統(tǒng)計的顏色值分布結(jié)果，cpuColorBuffer是cpu統(tǒng)計的顏色值分布結(jié)果射窒。
理論上結(jié)果應(yīng)該接近藏杖，但實際上printf出來的差異非常多。

        for (int i = 0; i < 3; ++i) {
            for (int j = 0; j < LY_CHANNEL_SIZE; ++j) {
                val[i] += buffer->channel[i][j];
                convertBuffer->channel[i][j] = val[i] * 1.0 * (LY_CHANNEL_SIZE - 1) / sum;
                
                // 對比CPU和GPU處理的結(jié)果
                if (buffer->channel[i][j] != strongSelf->cpuColorBuffer.channel[i][j]) {
                    // 如果不相同脉顿，則把對應(yīng)的結(jié)果輸出
                    printf("%d, %d, gpuBuffer:%u  cpuBuffer:%u \n", i, j, buffer->channel[i][j], strongSelf->cpuColorBuffer.channel[i][j]);
                }
            }
        }

問題分析：
通過檢查代碼蝌麸，先判定cpu統(tǒng)計的結(jié)果是正常。（cpu的處理過程就是正常的for循環(huán)弊予，不易出錯）
仔細觀察log的不同：
0, 1, gpuBuffer:763 cpuBuffer:762
結(jié)果很接近祥楣，但是有細微的差距。
我們知道gpu是浮點數(shù)的處理汉柒，而cpu是整數(shù)型處理误褪，浮點數(shù)到整數(shù)中間有精度的問題。
此時再看我們的shader碾褂，我們是以half來進行計算兽间，這樣統(tǒng)計出來的結(jié)果會有點誤差。

grayKernel(texture2d<half, access::read>  sourceTexture  [[texture(LYFragmentTextureIndexTextureSource)]],
           device LYColorBuffer &out [[buffer(LYKernelBufferIndexOutput)]],
           uint2                          grid         [[thread_position_in_grid]])

通過把精度從half改成float正塌，cpu和gpu的統(tǒng)計差異就只有3個：

0, 248, gpuBuffer:23215  cpuBuffer:22854
1, 74, gpuBuffer:23201  cpuBuffer:22840
2, 64, gpuBuffer:23336  cpuBuffer:22975

3嘀略、gpu渲染的圖片為白色

問題表現(xiàn)：
在gpu統(tǒng)計的結(jié)果與cpu接近的情況下，把映射buffer傳給fragment shader乓诽，最后進行一次顏色處理帜羊。
但是結(jié)果是白色的圖片，shader的代碼如下：

fragment float4
samplingShader(RasterizerData input [[stage_in]], // stage_in表示這個數(shù)據(jù)來自光柵化鸠天。（光柵化是頂點處理之后的步驟讼育，業(yè)務(wù)層無法修改）
               texture2d<float> colorTexture [[ texture(LYFragmentTextureIndexTextureSource) ]], // texture表明是紋理數(shù)據(jù)，LYFragmentTextureIndexTextureSource是索引
               device LYLocalBuffer &localBuffer [[buffer(LYFragmentBufferIndexConvert)]])
{
    constexpr sampler textureSampler (mag_filter::linear,
                                      min_filter::linear); // sampler是采樣器
    
    float4 colorSample = colorTexture.sample(textureSampler, input.textureCoordinate); // 得到紋理對應(yīng)位置的顏色
    int3 rgb = int3(colorSample.rgb);
    colorSample.rgb = float3(localBuffer.channel[0][rgb.r], localBuffer.channel[1][rgb.g], localBuffer.channel[2][rgb.b]);
    
    return colorSample;
}

問題分析：
我們先把colorSample.rgb = ...的這行代碼屏蔽稠集，發(fā)現(xiàn)渲染結(jié)果是正常的奶段，那么問題就出現(xiàn)在映射處理上面。
再通過Xcode的Capture GPU Frame工具剥纷，查看傳入的映射buffer數(shù)據(jù)痹籍，也是正常的數(shù)據(jù)。
那么問題可能出現(xiàn)int3 rgb的初始化晦鞋，或者從映射buffer讀取數(shù)據(jù)蹲缠。
觀察到int3 rgb = int3(colorSample.rgb)，是有一個float->int的操作悠垛，聯(lián)想到前面提到的歸一化處理吼砂，馬上明白：在這里的初始化時應(yīng)該乘以SIZE。

那么問題是否就此解決鼎文？不是的渔肩。

我們在進行顏色轉(zhuǎn)換的時候，float->int 需要乘以SIZE拇惋；
在獲取到映射buffer里面對應(yīng)顏色的值后周偎，仍需要做一次int->float的處理，除以SIZE撑帖；
如果下：

    float4 colorSample = colorTexture.sample(textureSampler, input.textureCoordinate); // 得到紋理對應(yīng)位置的顏色
    int3 rgb = int3(colorSample.rgb * size);
    colorSample.rgb = float3(localBuffer.channel[0][rgb.r], localBuffer.channel[1][rgb.g], localBuffer.channel[2][rgb.b]) / size;
    

4蓉坎、映射結(jié)果異常

問題表現(xiàn)：
問題如下，映射結(jié)果應(yīng)該是0~255的值胡嘿，但是通過Xcode看到最終的映射值遠超過255蛉艾，甚至接近255*2的數(shù)字。

問題分析：
下面是映射的算法

    int rgb[3][LY_CHANNEL_SIZE], sum = (int)(width * height);
    int val[3] = {0};
    for (int i = 0; i < 3; ++i) {
        for (int j = 0; j < LY_CHANNEL_SIZE; ++j) {
            val[i] += cpuColorBuffer.channel[i][j];
            rgb[i][j] = val[i] * 1.0 * (LY_CHANNEL_SIZE - 1) / sum;
        }
    }

sum是固定值，LY_CHANNEL_SIZE是常量值256勿侯，那么映射結(jié)果超過255的原因就是val[i]的統(tǒng)計結(jié)果太大拓瞪！
通過Xcode調(diào)試，確實如此：

那么助琐，會是什么原因?qū)е拢?/p>

在看到結(jié)果接近255的兩倍時祭埂，大概猜測可能是重復(fù)運算導(dǎo)致。
我們的均衡化處理是在MTKView的回調(diào)進行兵钮，如下：

- (void)drawInMTKView:(MTKView *)view {
    [self customDraw];
}

這里會回調(diào)多次蛆橡，從而導(dǎo)致多次執(zhí)行compute shader的顏色統(tǒng)計，這里可以引入isDrawing的臨時變量解決：

- (void)drawInMTKView:(MTKView *)view {
    if (!self.isDrawing) {
        self.isDrawing = YES;
        [self customDraw];
    }
}

但是掘譬，問題并沒有徹底解決：首次統(tǒng)計正常泰演，但是第二次處理的時候就會累積上一次的統(tǒng)計值。
如何對值進行清理葱轩？（這里實際上只處理一次也行睦焕，但是debug過程中需要通過Xcode的GPU Capture Frame工具進行查看，而這個工具需要多次渲染）
我們知道MTLBuffer是cpu酿箭、gpu都可以操作的buffer复亏，那么在cpu直接清除這個數(shù)據(jù)即可。
在 commandBuffer addCompletedHandler:^(){}的結(jié)束回調(diào)中缭嫡，使用memset(buffer, 0, strongSelf.colorBuffer.length)清理統(tǒng)計結(jié)果缔御。

5、映射結(jié)果最大值為256

問題表現(xiàn)：
在踩過上面的各種坑之后妇蛀，直方圖均衡化的效果也已經(jīng)展現(xiàn)耕突，但是仍有一點小問題：
映射結(jié)果buffer的數(shù)字范圍是0~256，而不是255评架。

問題分析：
根據(jù)直方圖均衡化的算法眷茁，我們知道是因為像素顏色值的統(tǒng)計，結(jié)果稍微偏大纵诞。
回顧Compute Shader的代碼：

kernel void
grayKernel(texture2d<float, access::read>  sourceTexture  [[texture(LYFragmentTextureIndexTextureSource)]],
           device LYColorBuffer &out [[buffer(LYKernelBufferIndexOutput)]],
           uint2                          grid         [[thread_position_in_grid]]) {
    // 邊界保護
    if(grid.x <= sourceTexture.get_width() && grid.y <= sourceTexture.get_height())
    {
        float4 color  = sourceTexture.read(grid); // 初始顏色
        int3 rgb = int3(color.rgb * size); //
        atomic_fetch_add_explicit(&out.channel[0][rgb.r], 1, memory_order_relaxed);
        atomic_fetch_add_explicit(&out.channel[1][rgb.g], 1, memory_order_relaxed);
        atomic_fetch_add_explicit(&out.channel[2][rgb.b], 1, memory_order_relaxed);
    }
}

Metal的Compute Shader是按每組網(wǎng)格進行處理上祈，那么可能會出現(xiàn)邊界超過圖像的情況，所以添加了邊界保護浙芙。
但是登刺，這里存在誤判的情況：邊界判斷不應(yīng)該是<=，而是<嗡呼。
因為這個原因纸俭，會導(dǎo)致統(tǒng)計的結(jié)果偏大，最終出現(xiàn)256的情況南窗。

在修復(fù)完這個問題后揍很，Demo再無其他問題郎楼，GPU的處理結(jié)果也和CPU的處理結(jié)果完全一致！

總結(jié)

本文是在前文的Metal入門教程基礎(chǔ)上進行更復(fù)雜的嘗試窒悔，過程中也遇到較多問題呜袁，最終demo也順利完成，地址在Github蛉迹。

——-2018年8月31日 周五 23:51傅寡，踩著8月的尾巴更新這一篇文章放妈，九月是休假月北救。