OpenCL (Open Computing Language) 是由Khronos Group制定的开放式并行计算标准。它允许开发者利用异构计算平台上的多核CPU和GPU等处理器进行并行计算。OpenCL可以用于各种应用领域，包括科学计算、图形处理、机器学习等。
😁OpenCL的入门以及矢量相加的例子可以看我另一篇博客：
👉相关代码也可以直接访问git，这里贴上链接：
https://github.com/zly5/Parallel-Computing-Lab
https://github.com/zly5/Parallel-Computing-Lab
🚩先沾上目录：
目录
OpenCL矩阵乘法分析
OpenCL矩阵乘法的实现也是一种并行计算任务，允许在不同的计算设备上执行。以下是对OpenCL矩阵乘法实现的简要分析：
-
分解问题：
- 矩阵乘法的任务是将两个矩阵相乘以生成一个结果矩阵。
- 每个元素的计算是独立的，因此可以在并行计算设备上并行执行。
-
OpenCL Kernel：
- 在OpenCL中，矩阵乘法通常作为一个内核函数（kernel）来实现。
- 内核函数在OpenCL设备上并行执行，每个工作项（work item）计算一个结果矩阵元素。
- 全局工作大小和局部工作大小的设置用于确定如何分配工作项。
-
内存分配：
- 矩阵A、B和C通常存储在全局内存中，需要通过OpenCL API进行内存分配和数据传输。
- 主机（CPU）代码需要将数据传输到设备（如GPU）的全局内存中，以便内核函数可以访问。
- 数据传输是主机和设备之间的性能瓶颈，因此需要谨慎管理内存。
-
工作项分布：
- OpenCL中，工作项通常划分为工作组（work group）。
- 工作组是一组工作项，它们可以在局部内存中协同工作，并提高数据局部性。
- 工作组的划分和大小通常是根据设备特性和矩阵大小来优化的。
-
工作项计算：
- 每个工作项负责计算结果矩阵的一个元素。
- 工作项通过访问矩阵A和B的相应元素，执行乘法和求和操作，将结果写入矩阵C中。
- 工作组内的工作项可以通过局部内存共享中间计算结果，以减少全局内存访问。
-
优化：
- OpenCL矩阵乘法可以通过多种方式进行优化，如局部内存的使用、使用向量数据类型、调整工作组大小等。
- 优化旨在提高计算效率、减少内存带宽需求和减少访存延迟，以充分利用计算设备的性能。
总的来说，OpenCL矩阵乘法是一种可移植的并行计算问题，允许在不同的计算设备上执行，包括GPU、CPU和FPGA等。通过合理的内核设计和内存管理，可以加速矩阵乘法运算，提高计算效率。但需要深入了解OpenCL编程模型和设备特性，以实现最佳性能。同时，数据传输和内存访问是需要特别关注的性能瓶颈。
OpenCL实现矩阵乘法的细节
平台和设备的选择：
使用OpenCL API选择合适的平台和设备。
获取平台和设备信息，确定可以使用的计算资源。
内存分配和数据传输：
在主机（CPU）上分配内存用于存储矩阵A、B和C。
使用OpenCL API将数据传输到设备（如GPU）的全局内存中。
Kernel编写：
编写OpenCL内核函数，通常是一个矩阵乘法的函数。
内核函数定义如何计算一个结果矩阵元素。
内核函数可以使用get_global_id等函数获取工作项在全局工作空间中的ID。
Kernel编译和构建：
使用OpenCL API创建程序对象，将内核源代码加载到程序对象中。
调用API编译和构建程序，以生成可执行内核。
工作组和工作项的设置：
决定全局工作大小和局部工作大小。
设置工作组和工作项的数量，以适应矩阵的大小和设备的特性。
Kernel参数设置：
使用OpenCL API为内核函数设置参数，例如矩阵A、B、C的内存对象。
设置其他常量参数，如矩阵的维度等。
局部内存的使用：
在内核函数中，可以使用局部内存（local memory）来缓存部分矩阵，以减少全局内存访问。
适当使用局部内存可以提高内存访问的效率。
内核的执行：
使用OpenCL API调用内核函数，并传递全局工作大小和局部工作大小。
内核函数在设备上并行执行，每个工作项负责计算一个结果矩阵元素。
数据传输回主机：
使用OpenCL API将计算结果从设备传输回主机内存。
主机代码可以访问计算得到的结果矩阵C。
内存释放：
在程序结束时，使用OpenCL API释放设备上的内存对象，防止内存泄漏。
错误处理：
在代码中加入错误处理机制，以便检测和处理OpenCL API调用中的错误。
实验环境：
CPU型号:R7-5800H ，8核， 3.2GHz
GPU型号: NVIDIA GeForce GTX 3050，显存大小 4GB，CUDA核心：2560个，单精度性能为 9TFLOPS。
实验设计思路：
对于矩阵乘法，我们使用每一个GPU线程去计算矩阵A的一行和矩阵B的一列相点积的结果，并比较不同矩阵大小下的时间和加速比情况。
👉其程序设计大致可分为6个步骤，分别如下：
（1）获取计算平台(Platform)，查找支持OpenCL的硬件设备，并创建上下文(Context)； 函数clGetPlatformIDs可用来获取可用平台的数量和列表，函数clGetDeviceIDs用来获取OpenCL设备的数量和列表。
（2）创建命令队列(Command Queue)及包含了内核的程序(Program)对象，如果该程序是源代码，则还需进行在线编译； 需要使用函数clCreateCommandQueue创建命令队列，使用命令clCreateProgramWithSource创建程序对象，程序对象需要一个字符串指针，这个指针指向GPU计算核代码，使用命令clBuildProgram编译GPU并行程序。
（3） 创建程序执行过程中需要的存储对象(Buffer)，并初始化；
在GPU上利用clCreateBuffer函数开辟数据空间，需要设置GPU显存空间的读写属性和开辟空间的大小。利用函数clEnqueueWriteBuffer将数据从内存传输到显存。
（4） 创建内核对象并设置其所需参数；
使用命令clCreateKernel和核函数名称从程序对象中建立核函数对象。使用clSetKernelArg函数设置核程序参数。
（5） 设置内核的索引空间(NDRange)并执行内核，其中，NDRange通过全局尺寸(Global Size)和工作组尺寸(Work Group)来进行管理；
调用函数clEnqueueNDRangeKernel执行核代码。函数中要设置全局网格和局部网格的线程组织方式。然后调用clFinish函数以确保命令队列中的命令执行完毕。
（6） 将运行的结果拷贝回主机(Host)内存。
利用函数clEnqueueReadBuffer将计算结果从显存拷贝回主机内存。
OpenCL实验：
设置的BLOCK_SIZE为16，GRID_SIZE根据矩阵的实际大小进行计算。对于一个M*N的矩阵和N*M的矩阵相乘：
完整版代码如下：
#include <CL/OpenCL.h>
#include <iostream>
#include <fstream>
#include <sstream>
#include<time.h>
#include<stdio.h>
#include<stdlib.h>
using namespace std;
const int t = 1000;
const int heightA = t;
const int widthB = t;
const int midle = t;
//const int heightB = 3;
//一、 选择OpenCL平台并创建一个上下文
cl_context CreateContext()
{
cl_int errNum;
cl_uint numPlatforms;
cl_platform_id firstPlatformId;
cl_context context = NULL;
//选择可用的平台中的第一个
errNum = clGetPlatformIDs(1, &firstPlatformId, &numPlatforms);
if (errNum != CL_SUCCESS || numPlatforms <= 0)
{
std::cerr << "Failed to find any OpenCL platforms." << std::endl;
return NULL;
}
//创建一个OpenCL上下文环境
cl_context_properties contextProperties[] =
{
CL_CONTEXT_PLATFORM,
(cl_context_properties)firstPlatformId,
0
};
context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU,
NULL, NULL, &errNum);
return context;
}
//二、 创建设备并创建命令队列
cl_command_queue CreateCommandQueue(cl_context context, cl_device_id* device)
{
cl_int errNum;
cl_device_id* devices;
cl_command_queue commandQueue = NULL;
size_t deviceBufferSize = -1;
// 获取设备缓冲区大小
errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &deviceBufferSize);
if (deviceBufferSize <= 0)
{
std::cerr << "No devices available.";
return NULL;
}
// 为设备分配缓存空间
devices = new cl_device_id[deviceBufferSize / sizeof(cl_device_id)];
errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, deviceBufferSize, devices, NULL);
//选取可用设备中的第一个
commandQueue = clCreateCommandQueue(context, devices[0], 0, NULL);
*device = devices[0];
delete[] devices;
return commandQueue;
}
// 三、创建和构建程序对象
cl_program CreateProgram(cl_context context, cl_device_id device, const char* fileName)
{
cl_int errNum;
cl_program program;
std::ifstream kernelFile(fileName, std::ios::in);
if (!kernelFile.is_open())
{
std::cerr << "Failed to open file for reading: " << fileName << std::endl;
return NULL;
}
std::ostringstream oss;
oss << kernelFile.rdbuf();
std::string srcStdStr = oss.str();
const char* srcStr = srcStdStr.c_str();
program = clCreateProgramWithSource(context, 1,
(const char**)&srcStr,
NULL, NULL);
errNum = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
return program;
}
//创建和构建程序对象
bool CreateMemObjects(cl_context context, cl_mem memObjects[3],
int* a, int* b)
{
memObjects[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
sizeof(int) * midle * heightA, a, NULL);
memObjects[1] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
sizeof(int) * widthB * midle, b, NULL);
memObjects[2] = clCreateBuffer(context, CL_MEM_READ_WRITE,
sizeof(int) * widthB * heightA, NULL, NULL);
return true;
}
// 释放OpenCL资源
void Cleanup(cl_context context, cl_command_queue commandQueue,
cl_program program, cl_kernel kernel, cl_mem memObjects[3])
{
for (int i = 0; i < 3; i++)
{
if (memObjects[i] != 0)
clReleaseMemObject(memObjects[i]);
}
if (commandQueue != 0)
clReleaseCommandQueue(commandQueue);
if (kernel != 0)
clReleaseKernel(kernel);
if (program != 0)
clReleaseProgram(program);
if (context != 0)
clReleaseContext(context);
}
int main(int argc, char** argv)
{
cl_context context = 0;
cl_command_queue commandQueue = 0;
cl_program program = 0;
cl_device_id device = 0;
cl_kernel kernel = 0;
cl_mem memObjects[3] = { 0, 0, 0 };
cl_int errNum;
cl_event events[1];
clock_t t1, t2, t3;
const char* filename = "./a.cl";
// 一、选择OpenCL平台并创建一个上下文
context = CreateContext();
// 二、 创建设备并创建命令队列
commandQueue = CreateCommandQueue(context, &device);
//三、创建和构建程序对象
program = CreateProgram(context, device, filename);
// 四、 创建OpenCL内核并分配内存空间
kernel = clCreateKernel(program, "hello_kernel", NULL);
//创建要处理的数据
int* a = NULL; // 输入数组
int* b = NULL; // 输入数组
int* result = NULL; // 输出数组
// 数组的大小
const int elementsA = heightA * midle;
const int elementsB = midle * widthB;
const int elementsC = heightA * widthB;
// 计算内存大小
size_t datasizeA = sizeof(float) * elementsA;
size_t datasizeB = sizeof(float) * elementsB;
size_t datasizeC = sizeof(float) * elementsC;
// 分配内存空间
a = (int*)malloc(datasizeA);
b = (int*)malloc(datasizeB);
result = (int*)malloc(datasizeC);
for (int i = 0; i < heightA; i++)
{
for (int j = 0; j < midle; j++)
{
a[i * midle + j] = 2;//10.0f * ((int) rand() / (int) RAND_MAX);
}
}
for (int k = 0; k < midle; k++)
{
for (int m = 0; m < widthB; m++)
{
b[k * widthB + m] = 3;//10.0f * ((int) rand() / (int) RAND_MAX);
}
}
t1 = clock(); //mach_absolute_time();
//cpu串行处理代码
for (int l = 0; l < heightA; l++) {
for (int n = 0; n < widthB; n++) {
for (int q = 0; q < midle; q++) {
result[l * widthB + n] += a[l * midle + q] * b[q * widthB + n];
}
//std::cout<<"r = "<<result[l*widthB+n]<<std::endl;
}
}
t2 = clock(); //mach_absolute_time();
//创建内存对象
if (!CreateMemObjects(context, memObjects, a, b))
{
Cleanup(context, commandQueue, program, kernel, memObjects);
return 1;
}
// 五、 设置内核数据并执行内核
errNum = clSetKernelArg(kernel, 0, sizeof(cl_mem), &memObjects[0]);
errNum = clSetKernelArg(kernel, 1, sizeof(cl_mem), &memObjects[1]);
errNum = clSetKernelArg(kernel, 2, sizeof(cl_mem), &memObjects[2]);
errNum = clSetKernelArg(kernel, 3, sizeof(int), &heightA);
errNum = clSetKernelArg(kernel, 4, sizeof(int), &widthB);
errNum = clSetKernelArg(kernel, 5, sizeof(int), &midle);
size_t globalWorkSize[2];
globalWorkSize[0] = heightA;
globalWorkSize[1] = widthB;
// size_t localWorkSize[2] = { 1,1 };
errNum = clEnqueueNDRangeKernel(commandQueue, kernel, 2, NULL,
globalWorkSize, NULL,
0, NULL, &events[0]);
// 六、 读取执行结果并释放OpenCL资源
errNum = clEnqueueReadBuffer(commandQueue, memObjects[2], CL_TRUE,
0, widthB * heightA * sizeof(int), result,
0, NULL, NULL);
// for(int p=0;p<20;p++){
// cout<<"new ="<<result[p];
// }
errNum = clWaitForEvents(1, &events[0]);
t3 = clock(); //mach_absolute_time();
errNum = clReleaseEvent(events[0]);
printf("cpu t = %.8f\n", ((double)t2 - (double)t1) / CLOCKS_PER_SEC);
printf("gpu t = %.8f \n", ((double)t3 - (double)t2)/ CLOCKS_PER_SEC);
std::cout << std::endl;
std::cout << "Executed program succesfully." << std::endl;
Cleanup(context, commandQueue, program, kernel, memObjects);
return 0;
}其中a.cl文件如下：
__kernel void hello_kernel(__global const int *a,
__global const int *b,
__global int *result_matrix,int result_matrix_row,
int result_matrix_col,int compute_size)
{
int row = get_global_id(0);
int col = get_global_id(1);
int sum = 0;
for(int i=0;i<compute_size;i++)
{
sum += a[row*compute_size+i] * b[i*result_matrix_col+col];
}
result_matrix[row*result_matrix_col+col] = sum;
}分别测试了以下不同大小的矩阵在CPU和GPU上的计算结果，具体如下表1所示。
表1 OpenCL并行矩阵乘法实验结果
| 矩阵大小 | CPU运行时间（s） | GPU运行时间（s） | 加速比 |
| 300*300 × 300*300 | 0.062 | 0.002 | 31 |
| 400*400 × 400*400 | 0.146 | 0.004 | 36.5 |
| 500*500 × 500*500 | 0.288 | 0.006 | 48 |
| 600*600 × 600*600 | 0.538 | 0.009 | 59.77778 |
| 700*700 × 700*700 | 0.916 | 0.012 | 76.33333 |
| 800*800 × 800*800 | 1.225 | 0.039 | 31.41026 |
| 900*900 × 900*900 | 1.999 | 0.024 | 83.29167 |
加速比：是同一个任务在单处理器系统和并行处理器系统中运行消耗的时间的比率，用来衡量并行系统或程序并行化的性能和效果。其定义如下式所示：
其中，Sp是加速比，T1是单处理器下的运行时间，Tp是在有p个处理器并行系统中的运行时间。当Sp＝p时，此加速比被称为线性加速比。
❤️总结：
从OpenL的实验结果可以看到，并行OpenCL程序实现的矩阵乘法相比于串行的串行有了很大的速度提升，而已我们观察到，对于不同阶数的矩阵计算来说，较小的矩阵加速比往往比较大矩阵的加速比要低，这可能是因为计算任务小，GPU资源利用率不高，相较于计算来说，可能内核启动、分配内存所占时间的比重更大。
扫一扫
9866
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
