本公眾號內容均為本號轉發，已儘可能註明出處。因未能核實來源或轉發內容圖片有權利瑕疵的，請及時聯繫本號，本號會第一時間進行修改或刪除。 QQ : 3442093904
前言
Metal入門教程總結
正文
核心思路
首先，我們用直方圖來表示一張圖像：橫坐標代表的是顏色值，縱坐標代表的是該顏色值在圖像中出現次數。
如圖，對於某些圖像，可能出現顏色值集中分布在某個區間的情況。
直方圖均衡化(Histogram Equalization) ，指的是對圖像的顏色值進行重新分配，使得顏色值的分布更加均勻。
本文用compute shader對圖像的顏色值進行統計，然後計算得出映射關係，由fragment shader進行顏色映射處理。
效果展示
具體步驟
1、Metal的render管道、compute管道配置；
同前文，不再贅述，詳見Metal入門教程總結。
2、CPU進行直方圖均衡化處理；
2.1 把UIImage轉成Bytes；
2.2 顏色統計；
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 映射關係；
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];
}
}
3 GPU進行直方圖均衡化處理；
kernel voidgrayKernel(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]的顏色值
// 顏色統計，每個像素點計一次
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);
}
}</float, access::read>
atomic_fetch_add_explicit是用於在多線程進行數據操作，具體的函數解釋見這裡。
compute shader回調後，根據GPU統計的顏色分布結果，求出映射關係；
LYLocalBuffer *buffer = (LYLocalBuffer *)strongSelf.colorBuffer.contents; // GPU統計的結果
LYLocalBuffer *convertBuffer = self.convertBuffer.contents; // 顏色轉換的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]累計出現的總次數
convertBuffer->channel[i][j] = val[i] * 1.0 * (LY_CHANNEL_SIZE - 1) / sum;
// 對比CPU和GPU處理的結果
if (buffer->channel[i][j] != strongSelf->cpuColorBuffer.channel[i][j]) { // 如果不相同，則把對應的結果輸出
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);
fragment float4samplingShader(RasterizerData input [[stage_in]], // stage_in表示這個數據來自光柵化。（光柵化是頂點處理之後的步驟，業務層無法修改）
texture2d<float> colorTexture [[ texture(LYFragmentTextureIndexSource) ]], // texture表明是紋理數據，LYFragmentTextureIndexSource是索引
device LYLocalBuffer &convertBuffer [[buffer(LYFragmentBufferIndexConvert)]]) // 轉換的buffer{ constexpr sampler textureSampler (mag_filter::linear, min_filter::linear); // sampler是採樣器
float4 colorSample = colorTexture.sample(textureSampler, input.textureCoordinate); // 得到紋理對應位置的顏色
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; // 返回的值也要經過歸一化處理
return colorSample;
}</float>
遇到的問題
1、統計結果集中在頭部
問題表現：
統計結果異常，集中在前面兩個值。
如下，green通道的顏色集中在r[0]和r[1]上：
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的值；（歸一化）
統計的值全部是在前面，是因為沒有乘以255！
先用CPU實現了直方圖均衡化，在實現shader的時候，參考CPU的代碼實現，犯了這個錯誤。
2、cpu和gpu統計結果相差較多
問題表現：
如下代碼，buffer是gpu統計的顏色值分布結果，cpuColorBuffer是cpu統計的顏色值分布結果。
理論上結果應該接近，但實際上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處理的結果
if (buffer->channel[i][j] != strongSelf->cpuColorBuffer.channel[i][j]) { // 如果不相同，則把對應的結果輸出
printf("%d, %d, gpuBuffer:%u cpuBuffer:%u \n", i, j, buffer->channel[i][j], strongSelf->cpuColorBuffer.channel[i][j]);
}
}
}
問題分析：
通過檢查代碼，先判定cpu統計的結果是正常。（cpu的處理過程就是正常的for循環，不易出錯）
仔細觀察log的不同：
0, 1, gpuBuffer:763 cpuBuffer:762
結果很接近，但是有細微的差距。
我們知道gpu是浮點數的處理，而cpu是整數型處理，浮點數到整數中間有精度的問題。
此時再看我們的shader，我們是以half來進行計算，這樣統計出來的結果會有點誤差。
grayKernel(texture2d<half, access::read> sourceTexture [[texture(LYFragmentTextureIndexTextureSource)]],
device LYColorBuffer &out [[buffer(LYKernelBufferIndexOutput)]],
uint2 grid [[thread_position_in_grid]])</half, access::read>
通過把精度從half改成float，cpu和gpu的統計差異就只有3個：
0, 248, gpuBuffer:23215 cpuBuffer:22854
1, 74, gpuBuffer:23201 cpuBuffer:22840
2, 64, gpuBuffer:23336 cpuBuffer:22975
3、gpu渲染的圖片為白色
問題表現：
在gpu統計的結果與cpu接近的情況下，把映射buffer傳給fragment shader，最後進行一次顏色處理。
但是結果是白色的圖片，shader的代碼如下：
fragment float4samplingShader(RasterizerData input [[stage_in]], // stage_in表示這個數據來自光柵化。（光柵化是頂點處理之後的步驟，業務層無法修改）
texture2d<float> colorTexture [[ texture(LYFragmentTextureIndexTextureSource) ]], // texture表明是紋理數據，LYFragmentTextureIndexTextureSource是索引
device LYLocalBuffer &localBuffer [[buffer(LYFragmentBufferIndexConvert)]]){ constexpr sampler textureSampler (mag_filter::linear,
min_filter::linear); // sampler是採樣器
float4 colorSample = colorTexture.sample(textureSampler, input.textureCoordinate); // 得到紋理對應位置的顏色
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;
}</float>
問題分析：
我們先把colorSample.rgb = ...的這行代碼屏蔽，發現渲染結果是正常的，那麼問題就出現在映射處理上面。
再通過Xcode的Capture GPU Frame工具，查看傳入的映射buffer數據，也是正常的數據。
那麼問題可能出現int3 rgb的初始化，或者從映射buffer讀取數據。
觀察到int3 rgb = int3(colorSample.rgb)，是有一個float->int的操作，聯想到前面提到的歸一化處理，馬上明白：在這裡的初始化時應該乘以SIZE。
那麼問題是否就此解決？不是的。
我們在進行顏色轉換的時候，float->int 需要乘以SIZE；
在獲取到映射buffer裡面對應顏色的值後，仍需要做一次int->float的處理，除以SIZE；
如果下：
float4 colorSample = colorTexture.sample(textureSampler, input.textureCoordinate); // 得到紋理對應位置的顏色
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、映射結果異常
問題表現：
問題如下，映射結果應該是0~255的值，但是通過Xcode看到最終的映射值遠超過255，甚至接近255*2的數字。
問題分析：
下面是映射的算法
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，那麼映射結果超過255的原因就是val[i]的統計結果太大！
通過Xcode調試，確實如此：
那麼，會是什麼原因導致？
在看到結果接近255的兩倍時，大概猜測可能是重複運算導致。
我們的均衡化處理是在MTKView的回調進行，如下：
- (void)drawInMTKView:(MTKView *)view {
[self customDraw];
}
這裡會回調多次，從而導致多次執行compute shader的顏色統計，這裡可以引入isDrawing的臨時變量解決：
- (void)drawInMTKView:(MTKView *)view {
if (!self.isDrawing) {
self.isDrawing = YES;
[self customDraw];
}
}
但是，問題並沒有徹底解決：首次統計正常，但是第二次處理的時候就會累積上一次的統計值。
如何對值進行清理？（這裡實際上只處理一次也行，但是debug過程中需要通過Xcode的GPU Capture Frame工具進行查看，而這個工具需要多次渲染）
我們知道MTLBuffer是cpu、gpu都可以操作的buffer，那麼在cpu直接清除這個數據即可。
在 commandBuffer addCompletedHandler:^(){}的結束回調中，使用memset(buffer, 0, strongSelf.colorBuffer.length)清理統計結果。
5、映射結果最大值為256
問題表現：
在踩過上面的各種坑之後，直方圖均衡化的效果也已經展現，但是仍有一點小問題：
映射結果buffer的數字範圍是0~256，而不是255。
問題分析：
根據直方圖均衡化的算法，我們知道是因為像素顏色值的統計，結果稍微偏大。
回顧Compute Shader的代碼：
kernel voidgrayKernel(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);
}
}</float, access::read>
Metal的Compute Shader是按每組網格進行處理，那麼可能會出現邊界超過圖像的情況，所以添加了邊界保護。
但是，這裡存在誤判的情況：邊界判斷不應該是<=，而是<。
因為這個原因，會導致統計的結果偏大，最終出現256的情況。
在修復完這個問題後，Demo再無其他問題，GPU的處理結果也和CPU的處理結果完全一致！
總結
本文是在前文的Metal入門教程基礎上進行更複雜的嘗試，過程中也遇到較多問題，最終demo也順利完成，地址在Github。
作者：落影loyinglin
連結：https://www.jianshu.com/p/1c8e814edab4