Metal入门学习:GPU并行计算大数组相加
一、编程指南PDF下载链接(中英文档)
-
1、Metal编程指南PDF链接
https://github.com/dennie-lee/ios_tech_record/raw/main/Metal学习PDF/Metal 编程指南.pdf -
2、Metal着色语言(Metal Shader Language:简称MSL)编程指南PDF链接
https://github.com/dennie-lee/ios_tech_record/raw/main/Metal学习PDF/Metal 着色语言指南.pdf -
3、补充:官网API文档链接
https://developer.apple.com/documentation/metal/performing_calculations_on_a_gpu
二、内容前述
本文章通过元素个数相同的两个数组对应位置相加来切入Metal(GPU)的并行计算着色函数(本篇文章未涉及渲染函数:顶点着色函数和片元着色函数,会另起一篇文章介绍)。在此示例中,可以了解所有 Metal 应用程序中使用的基本任务。 您将看到如何将用 C 编写的简单函数转换为Metal着色语言 (Metal Shader Language :MSL),以便它可以在 GPU 上运行。通过创建管道,准备MSL函数在其上运行,并创建GPU可访问的数据对象。要针对您的数据执行管道,创建命令缓冲区,将命令写入其中,然后将缓冲区提交到命令队列,Metal将命令发送到GPU执行。
(下图是官网并行计算着色函数流程原理图:)
三、C语言和MSL语言对两个数组相加的函数对比
- 1、C语言函数:两个数组相加
void add_arrays(const float* inA,const float* inB,float* result,int length){for (int index = 0; index < length ; index++){result[index] = inA[index] + inB[index];}
}
- 2、MSL语言函数:两个数组相加
kernel void add_arrays(device const float* inA,device const float* inB,device float* result,uint index [[thread_position_in_grid]]){result[index] = inA[index] + inB[index];
}
关键词解释说明(在MSL编程指南PDF中可查看,点击上面的链接下载即可)
[[thread_position_in_grid]]文档中也有,更加详细的解析说明可查看这篇文章:https://juejin.cn/post/7085633906501746724,文章中还有延伸说明解释[[threadgroup_position_in_grid]]和[[threads_per_threadgroup]]
四、关键代码段解析
- 1、初始化、加载扩展名为.metal的文件、创建管道状态对象加载并行计算着色函数
//_mDevice = MTLCreateSystemDefaultDevice();
// Load the shader files with a .metal file extension in the project
id<MTLLibrary> newDefaultLibrary = [_mDevice newDefaultLibrary];
if (newDefaultLibrary == nil){NSLog(@"Failed to find the default library");return nil;
}id<MTLFunction> newFunction = [newDefaultLibrary newFunctionWithName:@"add_arrays"];
if (newFunction == nil){NSLog(@"Failed to find the adder function");return nil;
}NSError *error;
// Create a compute pipeline state object.
//根据扩展名为.metal文件中kernel定义的函数创建计算管道(项目Add文件)
_mAddFunctionPSO = [_mDevice newComputePipelineStateWithFunction:newFunction error:&error];
if (_mAddFunctionPSO == nil){// If the Metal API validation is enabled, you can find out more information about what// went wrong. (Metal API validation is enabled by default when a debug build is run// from Xcode)NSLog(@"Failed to create pipeline state object,error : %@)",error);return nil;
}// 指令队列
_mCommandQueue = [_mDevice newCommandQueue];
if (_mCommandQueue == nil){NSLog(@"Failed to find command queue");return nil;
}
- 2、初始化数组数据
//初始化数组数据
- (void)prepareData{_mBufferA = [_mDevice newBufferWithLength:bufferSize options:MTLResourceStorageModeShared];_mBufferB = [_mDevice newBufferWithLength:bufferSize options:MTLResourceStorageModeShared];_mBufferResult = [_mDevice newBufferWithLength:bufferSize options:MTLResourceStorageModeShared];[self generateRandomFloatData:_mBufferA];[self generateRandomFloatData:_mBufferB];
}- (void)generateRandomFloatData:(id<MTLBuffer>)buffer{float *dataPtr = buffer.contents;for (int index = 0; index < arrayLength; index++){dataPtr[index] = (float)rand() / (float)RAND_MAX;}
}
- 3、指令参数添加,提交GPU计算
- (void)sendComputeCommand{//create a command buffer to hold commands//创建指令缓存冲区id<MTLCommandBuffer> commandBuffer = [_mCommandQueue commandBuffer];assert(commandBuffer != nil);//开始进行指令添加参数id<MTLComputeCommandEncoder> computeEncoder = [commandBuffer computeCommandEncoder];assert(computeEncoder != nil);NSTimeInterval startTimeInterval = [[[NSDate alloc] init] timeIntervalSince1970];[self encodeAdderCommand:computeEncoder];//添加参数完毕,[computeEncoder endEncoding];//提交执行指令[commandBuffer commit];//等待计算完毕[commandBuffer waitUntilCompleted];NSTimeInterval endTimeInterval = [[[NSDate alloc] init] timeIntervalSince1970];NSLog(@"长度为:%u 的两个数组相加(Metal方式)花费的时间:%f",arrayLength,(endTimeInterval - startTimeInterval));//验证计算结果[self verifyResults];
}- (void)encodeAdderCommand:(id<MTLComputeCommandEncoder>)computeEncoder{//encode pipeline state object and its parameters[computeEncoder setComputePipelineState:_mAddFunctionPSO];//此处的atIndex为0,与MSL函数中参数对应顺序对应(也可以设置buffer(0),此处可先忽略buffer,详细信息可查看着色语言编程指南PDF)[computeEncoder setBuffer:_mBufferA offset:0 atIndex:0];//此处的atIndex为1,与MSL函数中参数对应顺序对应(也可以设置buffer(1),此处可先忽略buffer,详细信息可查看着色语言编程指南PDF)[computeEncoder setBuffer:_mBufferB offset:0 atIndex:1];//此处的atIndex为1,与MSL函数中参数对应顺序对应(也可以设置buffer(2),此处可先忽略buffer,详细信息可查看着色语言编程指南PDF)[computeEncoder setBuffer:_mBufferResult offset:0 atIndex:2];//全部线程数量,此处对应[[thread_position_in_grid]],理解此处请看第三.2对应关键词解释的文章MTLSize gridSize = MTLSizeMake(arrayLength, 1, 1);NSUInteger threadGroupSize = _mAddFunctionPSO.maxTotalThreadsPerThreadgroup;if (threadGroupSize > arrayLength){threadGroupSize = arrayLength;}//线程组大小MTLSize threadGroupsize = MTLSizeMake(threadGroupSize, 1, 1);[computeEncoder dispatchThreads:gridSize threadsPerThreadgroup:threadGroupsize];
}
五、完整代码
例子:github链接:https://github.com/dennie-lee/MetalArrayAddDemo