《Mali OpenCL SDK v1.1.0》教程样例之一“Hello World”

2024-03-08 07:18
文章标签 sdk 教程 hello world 样例 opencl v1.1 mali

本文主要是介绍《Mali OpenCL SDK v1.1.0》教程样例之一“Hello World”，希望对大家解决编程问题提供一定的参考价值，需要的开发者们随着小编来一起学习吧！

1、算法简述


　　实现矩阵相加：Cn = An + Bn。这个例子虽然很简单，但是由于矩阵元素之间相互独立，每个元素可以非常容易地进行并行计算，可以非常理想地在OpenCL中实现。



2. C/C++实现

　　

[cpp] view plain copy print ? 在CODE上查看代码片 派生到我的代码片
  1. /* 
  2.  * This confidential and proprietary software may be used only as 
  3.  * authorised by a licensing agreement from ARM Limited 
  4.  *    (C) COPYRIGHT 2013 ARM Limited 
  5.  *        ALL RIGHTS RESERVED 
  6.  * The entire notice above must be reproduced on all authorised 
  7.  * copies and copies may only be made to the extent permitted 
  8.  * by a licensing agreement from ARM Limited. 
  9.  */  
  10.   
  11. #include <iostream>  
  12.   
  13. using namespace std;  
  14.   
  15. /** 
  16.  * \brief Basic integer array addition implemented in C/C++. 
  17.  * \details A sample which shows how to add two integer arrays and store the result in a third array. 
  18.  *          No OpenCL code is used in this sample, only standard C/C++. The code executes only on the CPU. 
  19.  * \return The exit code of the application, non-zero if a problem occurred. 
  20.  */  
  21. int main(void)  
  22. {  
  23.     /* [Setup memory] */  
  24.     /* Number of elements in the arrays of input and output data. */  
  25.     int arraySize = 1000000;  
  26.   
  27.     /* Arrays to hold the input and output data. */  
  28.     int* inputA = new int[arraySize];  
  29.     int* inputB = new int[arraySize];  
  30.     int* output = new int[arraySize];  
  31.     /* [Setup memory] */  
  32.   
  33.     /* Fill the arrays with data. */  
  34.     for (int i = 0; i < arraySize; i++)  
  35.     {  
  36.         inputA[i] = i;  
  37.         inputB[i] = i;  
  38.     }  
  39.   
  40.     /* [C/C++ Implementation] */  
  41.     for (int i = 0; i < arraySize; i++)  
  42.     {  
  43.         output[i] = inputA[i] + inputB[i];  
  44.     }  
  45.     /* [C/C++ Implementation] */  
  46.   
  47.     /* Uncomment the following block to print results. */  
  48.     /* 
  49.     for (int i = 0; i < arraySize; i++) 
  50.     { 
  51.         cout << "i = " << i << ", output = " <<  output[i] << "\n"; 
  52.     } 
  53.     */  
  54.   
  55.     delete[] inputA;  
  56.     delete[] inputB;  
  57.     delete[] output;  
  58. }  


3 Open基本实现


3.1 内核代码实现


　　内核代码的实现如下，其中指针的修饰符restrict是C99中的关键字，只用于限定指针。该关键字用于告知编译器，所有修改该指针所指向内容的操作全部都是基于该指针的，即不存在其它进行修改操作的途径；这样的后果是帮助编译器进行更好的代码优化，生成更有效率的汇编代码。

[cpp] view plain copy print ? 在CODE上查看代码片 派生到我的代码片
  1. /* 
  2.  * This confidential and proprietary software may be used only as 
  3.  * authorised by a licensing agreement from ARM Limited 
  4.  *    (C) COPYRIGHT 2013 ARM Limited 
  5.  *        ALL RIGHTS RESERVED 
  6.  * The entire notice above must be reproduced on all authorised 
  7.  * copies and copies may only be made to the extent permitted 
  8.  * by a licensing agreement from ARM Limited. 
  9.  */  
  10.   
  11. /** 
  12.  * \brief Hello World kernel function. 
  13.  * \param[in] inputA First input array. 
  14.  * \param[in] inputB Second input array. 
  15.  * \param[out] output Output array. 
  16.  */  
  17. /* [OpenCL Implementation] */  
  18. __kernel void hello_world_opencl(__global int* restrict inputA,  
  19.                                  __global int* restrict inputB,  
  20.                                  __global int* restrict output)  
  21. {  
  22.     /* 
  23.      * Set i to be the ID of the kernel instance. 
  24.      * If the global work size (set by clEnqueueNDRangeKernel) is n, 
  25.      * then n kernels will be run and i will be in the range [0, n - 1]. 
  26.      */  
  27.     int i = get_global_id(0);  
  28.   
  29.     /* Use i as an index into the three arrays. */  
  30.     output[i] = inputA[i] + inputB[i];  
  31. }  
  32. /* [OpenCL Implementation] */  

3.2 宿主机代码实现


　　内核代码中并没有循环语句，只计算一个矩阵元素的值，每一个实例获得一个独一无二的所以需要运行的内核实例数目等同于矩阵元素个数。

[cpp] view plain copy print ? 在CODE上查看代码片 派生到我的代码片
  1. /* 
  2.     * Each instance of our OpenCL kernel operates on a single element of each array so the number of 
  3.     * instances needed is the number of elements in the array. 
  4.     */  
  5.    size_t globalWorksize[1] = {arraySize};  
  6.    /* Enqueue the kernel */  
  7.    if (!checkSuccess(clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, globalWorksize, NULL, 0, NULL, &event)))  
  8.    {  
  9.        cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  10.        cerr << "Failed enqueuing the kernel. " << __FILE__ << ":"<< __LINE__ << endl;  
  11.        return 1;  
  12.    }  

　　因为我们并没有设置内核间的依赖性，OpenCL设备可以用并行的方式自由地运行内核实例。现在并行化上的唯一限制是设备的容量。在前面的代码运行之前，需要建立OpenCL，下面分别介绍与建立OpenCL相关的各项内容。


　　因为现在的操作是在GPU而不是CPU中，我们需要知道任何使用数据的位置。知道数据是在GPU内存空间还是CPU内存空间是非常重要的。在桌面系统中，GPU和CPU有它们自己的内存空间，被相对低速率的总线分开，这意味着在GPU和CPU之间共享数据是一个代价高昂的操作。在大多数带Mali-T600系列GPU的嵌入式系统中，GPU和CPU共享同一个内存，因此这使得以相对低的代价共享GPU和CPU之间内存成为可能。


　　由于这些系统的差异，OpenCL支持多种分配和共享设备间内存的方式。下面是一种共享设备间内存的方式，目的是减少从一个设备到另一个设备的内存拷贝(在一个共享内存系统中)。


a. 要求OpenCL设备分配内存


　　在C/C++实现中，我们使用数组来分配内存。

[cpp] view plain copy print ? 在CODE上查看代码片 派生到我的代码片
  1. /* Number of elements in the arrays of input and output data. */  
  2. int arraySize = 1000000;  
  3. /* Arrays to hold the input and output data. */  
  4. int* inputA = new int[arraySize];  
  5. int* inputB = new int[arraySize];  
  6. int* output = new int[arraySize];  
　 　在OpenCL中，我们使用内存缓冲区。内存缓冲区其实是一定大小的内存块。为了分配缓冲区，我们如下做：

[cpp] view plain copy print ? 在CODE上查看代码片 派生到我的代码片
  1. /* Number of elements in the arrays of input and output data. */  
  2. cl_int arraySize = 1000000;  
  3. /* The buffers are the size of the arrays. */  
  4. size_t bufferSize = arraySize * sizeof(cl_int);  
  5. /* 
  6.  * Ask the OpenCL implementation to allocate buffers for the data. 
  7.  * We ask the OpenCL implemenation to allocate memory rather than allocating 
  8.  * it on the CPU to avoid having to copy the data later. 
  9.  * The read/write flags relate to accesses to the memory from within the kernel. 
  10.  */  
  11. bool createMemoryObjectsSuccess = true;  
  12. memoryObjects[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, bufferSize, NULL, &errorNumber);  
  13. createMemoryObjectsSuccess &= checkSuccess(errorNumber);  
  14. memoryObjects[1] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, bufferSize, NULL, &errorNumber);  
  15. createMemoryObjectsSuccess &= checkSuccess(errorNumber);  
  16. memoryObjects[2] = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, bufferSize, NULL, &errorNumber);  
  17. createMemoryObjectsSuccess &= checkSuccess(errorNumber);  
  18. if (!createMemoryObjectsSuccess)  
  19. {  
  20.     cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  21.     cerr << "Failed to create OpenCL buffer. " << __FILE__ << ":"<< __LINE__ << endl;  
  22.     return 1;  
  23. }  
　 　尽管这看上去更加复杂，但其实这里只有三个OpenCL API调用。唯一的区别是这里我们检查错误(这是一个好的做法)，而C++中并不用做。

b. 映射内存到局部指针


　　现在内存已分配，但是只有OpenCL实现知道它的位置。为了访问CPU上的内存，我们把它们映射到一个指针。

[cpp] view plain copy print ? 在CODE上查看代码片 派生到我的代码片
  1. /* Map the memory buffers created by the OpenCL implementation to pointers so we can access them on the CPU. */  
  2. bool mapMemoryObjectsSuccess = true;  
  3. cl_int* inputA = (cl_int*)clEnqueueMapBuffer(commandQueue, memoryObjects[0], CL_TRUE, CL_MAP_WRITE, 0, bufferSize, 0, NULL, NULL, &errorNumber);  
  4. mapMemoryObjectsSuccess &= checkSuccess(errorNumber);  
  5. cl_int* inputB = (cl_int*)clEnqueueMapBuffer(commandQueue, memoryObjects[1], CL_TRUE, CL_MAP_WRITE, 0, bufferSize, 0, NULL, NULL, &errorNumber);  
  6. mapMemoryObjectsSuccess &= checkSuccess(errorNumber);  
  7. if (!mapMemoryObjectsSuccess)  
  8. {  
  9.    cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  10.    cerr << "Failed to map buffer. " << __FILE__ << ":"<< __LINE__ << endl;  
  11.    return 1;  
  12. }  

　　现在这些指针可以想普通的C/C++指针那样使用了。


c. 在CPU上初始化数据


　　因为我们已有了指向内存的指针，这一步与在CPU上一样。

[cpp] view plain copy print ? 在CODE上查看代码片 派生到我的代码片
  1. for (int i = 0; i < arraySize; i++)  
  2. {  
  3.    inputA[i] = i;  
  4.    inputB[i] = i;  
  5. }  

d. 取消映射缓冲区


　　为了使OpenCL设备使用缓冲区，我们必须把它们在CPU上的映射取消。

[cpp] view plain copy print ? 在CODE上查看代码片 派生到我的代码片
  1. /* 
  2.  * Unmap the memory objects as we have finished using them from the CPU side. 
  3.  * We unmap the memory because otherwise: 
  4.  * - reads and writes to that memory from inside a kernel on the OpenCL side are undefined. 
  5.  * - the OpenCL implementation cannot free the memory when it is finished. 
  6.  */  
  7. if (!checkSuccess(clEnqueueUnmapMemObject(commandQueue, memoryObjects[0], inputA, 0, NULL, NULL)))  
  8. {  
  9.    cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  10.    cerr << "Unmapping memory objects failed " << __FILE__ << ":"<< __LINE__ << endl;  
  11.    return 1;  
  12. }  
  13. if (!checkSuccess(clEnqueueUnmapMemObject(commandQueue, memoryObjects[1], inputB, 0, NULL, NULL)))  
  14. {  
  15.    cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  16.    cerr << "Unmapping memory objects failed " << __FILE__ << ":"<< __LINE__ << endl;  
  17.    return 1;  
  18. }  

e. 映射数据到内核


　　在我们调度内核运行之前，我们必须告诉内核哪些数据作为输入使用。这里，我们映射内存对象到OpenCL内核函数的参数中。

[cpp] view plain copy print ? 在CODE上查看代码片 派生到我的代码片
  1. bool setKernelArgumentsSuccess = true;  
  2. setKernelArgumentsSuccess &= checkSuccess(clSetKernelArg(kernel, 0, sizeof(cl_mem), &memoryObjects[0]));  
  3. setKernelArgumentsSuccess &= checkSuccess(clSetKernelArg(kernel, 1, sizeof(cl_mem), &memoryObjects[1]));  
  4. setKernelArgumentsSuccess &= checkSuccess(clSetKernelArg(kernel, 2, sizeof(cl_mem), &memoryObjects[2]));  
  5. if (!setKernelArgumentsSuccess)  
  6. {  
  7.     cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  8.     cerr << "Failed setting OpenCL kernel arguments. " << __FILE__ << ":"<< __LINE__ << endl;  
  9.     return 1;  
  10. }  

f. 运行内核


　　对于内核代码见前面，如何调度它则不作详述。


g. 获取运行结果


　　一旦计算结束，我们像映射输入缓冲区那样映射输出缓冲区。然后，我们就可以使用指针读取结果数据，然后取消缓冲区映射，就像前面那样。


　　基本实现的宿主机的完整代码如下：

[cpp] view plain copy print ? 在CODE上查看代码片 派生到我的代码片
  1. /* 
  2.  * This confidential and proprietary software may be used only as 
  3.  * authorised by a licensing agreement from ARM Limited 
  4.  *    (C) COPYRIGHT 2013 ARM Limited 
  5.  *        ALL RIGHTS RESERVED 
  6.  * The entire notice above must be reproduced on all authorised 
  7.  * copies and copies may only be made to the extent permitted 
  8.  * by a licensing agreement from ARM Limited. 
  9.  */  
  10.   
  11. #include "common.h"  
  12. #include "image.h"  
  13.   
  14. #include <CL/cl.h>  
  15. #include <iostream>  
  16.   
  17. using namespace std;  
  18.   
  19. /** 
  20.  * \brief Basic integer array addition implemented in OpenCL. 
  21.  * \details A sample which shows how to add two integer arrays and store the result in a third array. 
  22.  *          The main calculation code is in an OpenCL kernel which is executed on a GPU device. 
  23.  * \return The exit code of the application, non-zero if a problem occurred. 
  24.  */  
  25. int main(void)  
  26. {  
  27.     cl_context context = 0;  
  28.     cl_command_queue commandQueue = 0;  
  29.     cl_program program = 0;  
  30.     cl_device_id device = 0;  
  31.     cl_kernel kernel = 0;  
  32.     int numberOfMemoryObjects = 3;  
  33.     cl_mem memoryObjects[3] = {0, 0, 0};  
  34.     cl_int errorNumber;  
  35.   
  36.     if (!createContext(&context))  
  37.     {  
  38.         cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  39.         cerr << "Failed to create an OpenCL context. " << __FILE__ << ":"<< __LINE__ << endl;  
  40.         return 1;  
  41.     }  
  42.   
  43.     if (!createCommandQueue(context, &commandQueue, &device))  
  44.     {  
  45.         cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  46.         cerr << "Failed to create the OpenCL command queue. " << __FILE__ << ":"<< __LINE__ << endl;  
  47.         return 1;  
  48.     }  
  49.   
  50.     if (!createProgram(context, device, "assets/hello_world_opencl.cl", &program))  
  51.     {  
  52.         cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  53.         cerr << "Failed to create OpenCL program." << __FILE__ << ":"<< __LINE__ << endl;  
  54.         return 1;  
  55.     }  
  56.   
  57.     kernel = clCreateKernel(program, "hello_world_opencl", &errorNumber);  
  58.     if (!checkSuccess(errorNumber))  
  59.     {  
  60.         cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  61.         cerr << "Failed to create OpenCL kernel. " << __FILE__ << ":"<< __LINE__ << endl;  
  62.         return 1;  
  63.     }  
  64.   
  65.     /* [Setup memory] */  
  66.     /* Number of elements in the arrays of input and output data. */  
  67.     cl_int arraySize = 1000000;  
  68.   
  69.     /* The buffers are the size of the arrays. */  
  70.     size_t bufferSize = arraySize * sizeof(cl_int);  
  71.   
  72.     /* 
  73.      * Ask the OpenCL implementation to allocate buffers for the data. 
  74.      * We ask the OpenCL implemenation to allocate memory rather than allocating 
  75.      * it on the CPU to avoid having to copy the data later. 
  76.      * The read/write flags relate to accesses to the memory from within the kernel. 
  77.      */  
  78.     bool createMemoryObjectsSuccess = true;  
  79.   
  80.     memoryObjects[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, bufferSize, NULL, &errorNumber);  
  81.     createMemoryObjectsSuccess &= checkSuccess(errorNumber);  
  82.   
  83.     memoryObjects[1] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, bufferSize, NULL, &errorNumber);  
  84.     createMemoryObjectsSuccess &= checkSuccess(errorNumber);  
  85.   
  86.     memoryObjects[2] = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, bufferSize, NULL, &errorNumber);  
  87.     createMemoryObjectsSuccess &= checkSuccess(errorNumber);  
  88.   
  89.     if (!createMemoryObjectsSuccess)  
  90.     {  
  91.         cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  92.         cerr << "Failed to create OpenCL buffer. " << __FILE__ << ":"<< __LINE__ << endl;  
  93.         return 1;  
  94.     }  
  95.     /* [Setup memory] */  
  96.   
  97.     /* [Map the buffers to pointers] */  
  98.     /* Map the memory buffers created by the OpenCL implementation to pointers so we can access them on the CPU. */  
  99.     bool mapMemoryObjectsSuccess = true;  
  100.   
  101.     cl_int* inputA = (cl_int*)clEnqueueMapBuffer(commandQueue, memoryObjects[0], CL_TRUE, CL_MAP_WRITE, 0, bufferSize, 0, NULL, NULL, &errorNumber);  
  102.     mapMemoryObjectsSuccess &= checkSuccess(errorNumber);  
  103.   
  104.     cl_int* inputB = (cl_int*)clEnqueueMapBuffer(commandQueue, memoryObjects[1], CL_TRUE, CL_MAP_WRITE, 0, bufferSize, 0, NULL, NULL, &errorNumber);  
  105.     mapMemoryObjectsSuccess &= checkSuccess(errorNumber);  
  106.   
  107.     if (!mapMemoryObjectsSuccess)  
  108.     {  
  109.        cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  110.        cerr << "Failed to map buffer. " << __FILE__ << ":"<< __LINE__ << endl;  
  111.        return 1;  
  112.     }  
  113.     /* [Map the buffers to pointers] */  
  114.   
  115.     /* [Initialize the input data] */  
  116.     for (int i = 0; i < arraySize; i++)  
  117.     {  
  118.        inputA[i] = i;  
  119.        inputB[i] = i;  
  120.     }  
  121.     /* [Initialize the input data] */  
  122.   
  123.     /* [Un-map the buffers] */  
  124.     /* 
  125.      * Unmap the memory objects as we have finished using them from the CPU side. 
  126.      * We unmap the memory because otherwise: 
  127.      * - reads and writes to that memory from inside a kernel on the OpenCL side are undefined. 
  128.      * - the OpenCL implementation cannot free the memory when it is finished. 
  129.      */  
  130.     if (!checkSuccess(clEnqueueUnmapMemObject(commandQueue, memoryObjects[0], inputA, 0, NULL, NULL)))  
  131.     {  
  132.        cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  133.        cerr << "Unmapping memory objects failed " << __FILE__ << ":"<< __LINE__ << endl;  
  134.        return 1;  
  135.     }  
  136.   
  137.     if (!checkSuccess(clEnqueueUnmapMemObject(commandQueue, memoryObjects[1], inputB, 0, NULL, NULL)))  
  138.     {  
  139.        cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  140.        cerr << "Unmapping memory objects failed " << __FILE__ << ":"<< __LINE__ << endl;  
  141.        return 1;  
  142.     }  
  143.     /* [Un-map the buffers] */  
  144.   
  145.     /* [Set the kernel arguments] */  
  146.     bool setKernelArgumentsSuccess = true;  
  147.     setKernelArgumentsSuccess &= checkSuccess(clSetKernelArg(kernel, 0, sizeof(cl_mem), &memoryObjects[0]));  
  148.     setKernelArgumentsSuccess &= checkSuccess(clSetKernelArg(kernel, 1, sizeof(cl_mem), &memoryObjects[1]));  
  149.     setKernelArgumentsSuccess &= checkSuccess(clSetKernelArg(kernel, 2, sizeof(cl_mem), &memoryObjects[2]));  
  150.   
  151.     if (!setKernelArgumentsSuccess)  
  152.     {  
  153.         cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  154.         cerr << "Failed setting OpenCL kernel arguments. " << __FILE__ << ":"<< __LINE__ << endl;  
  155.         return 1;  
  156.     }  
  157.     /* [Set the kernel arguments] */  
  158.   
  159.     /* An event to associate with the Kernel. Allows us to retrieve profiling information later. */  
  160.     cl_event event = 0;  
  161.   
  162.     /* [Global work size] */  
  163.     /* 
  164.      * Each instance of our OpenCL kernel operates on a single element of each array so the number of 
  165.      * instances needed is the number of elements in the array. 
  166.      */  
  167.     size_t globalWorksize[1] = {arraySize};  
  168.     /* Enqueue the kernel */  
  169.     if (!checkSuccess(clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, globalWorksize, NULL, 0, NULL, &event)))  
  170.     {  
  171.         cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  172.         cerr << "Failed enqueuing the kernel. " << __FILE__ << ":"<< __LINE__ << endl;  
  173.         return 1;  
  174.     }  
  175.     /* [Global work size] */  
  176.   
  177.     /* Wait for kernel execution completion. */  
  178.     if (!checkSuccess(clFinish(commandQueue)))  
  179.     {  
  180.         cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  181.         cerr << "Failed waiting for kernel execution to finish. " << __FILE__ << ":"<< __LINE__ << endl;  
  182.         return 1;  
  183.     }  
  184.   
  185.     /* Print the profiling information for the event. */  
  186.     printProfilingInfo(event);  
  187.     /* Release the event object. */  
  188.     if (!checkSuccess(clReleaseEvent(event)))  
  189.     {  
  190.        cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  191.        cerr << "Failed releasing the event object. " << __FILE__ << ":"<< __LINE__ << endl;  
  192.        return 1;  
  193.     }  
  194.   
  195.     /* Get a pointer to the output data. */  
  196.     cl_int* output = (cl_int*)clEnqueueMapBuffer(commandQueue, memoryObjects[2], CL_TRUE, CL_MAP_READ, 0, bufferSize, 0, NULL, NULL, &errorNumber);  
  197.     if (!checkSuccess(errorNumber))  
  198.     {  
  199.        cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  200.        cerr << "Failed to map buffer. " << __FILE__ << ":"<< __LINE__ << endl;  
  201.        return 1;  
  202.     }  
  203.   
  204.     /* [Output the results] */  
  205.     /* Uncomment the following block to print results. */  
  206.     /* 
  207.     for (int i = 0; i < arraySize; i++) 
  208.     { 
  209.         cout << "i = " << i << ", output = " <<  output[i] << "\n"; 
  210.     } 
  211.     */  
  212.     /* [Output the results] */  
  213.   
  214.     /* Unmap the memory object as we are finished using them from the CPU side. */  
  215.     if (!checkSuccess(clEnqueueUnmapMemObject(commandQueue, memoryObjects[2], output, 0, NULL, NULL)))  
  216.     {  
  217.        cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  218.        cerr << "Unmapping memory objects failed " << __FILE__ << ":"<< __LINE__ << endl;  
  219.        return 1;  
  220.     }  
  221.   
  222.     /* Release OpenCL objects. */  
  223.     cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  224. }  


4 向量化你的OpenCL代码


4.1 向量基础


　　OpenCL设备可以通告它们为不同数据类型的首选向量宽度，你可以使用这个信息来选择一个内核。结果是，相当于该内核为你正在运行的平台做了优化。例如，一个设备可能仅有标量整数的硬件支持，而另一个设备则有宽度为4的整数向量的硬件支持。可以写两个版本的内核，一个用于标量，一个用于向量，在运行时选择正确的版本。

　　这里是一个在特定设备上询问首选整数向量宽度的例子。

[cpp] view plain copy print ? 在CODE上查看代码片 派生到我的代码片
  1. /* 
  2.  * Query the device to find out it's prefered integer vector width. 
  3.  * Although we are only printing the value here, it can be used to select between 
  4.  * different versions of a kernel. 
  5.  */  
  6. cl_uint integerVectorWidth;  
  7. clGetDeviceInfo(device, CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT, sizeof(cl_uint), &integerVectorWidth, NULL);  
  8. cout << "Prefered vector width for integers: " << integerVectorWidth << endl;  
对于其它OpenCL数据类型也是一样的。

　　每一个Mali T600系列GPU核最少有两个128位宽度的ALU(算数逻辑单元)，它们具有矢量计算能力。ALU中的绝大多数操作(例如，浮点加，浮点乘，整数加，整数乘)，可以以128位向量数据操作(例如，char16, short8, int4, float4)。使用前面讲述的询问方法来为你的数据类型决定使用正确的向量大小。

　　当使用Mali T600系列GPU时，我们推荐在任何可能的地方使用向量。


4.2 向量化代码


　　首先，修改内核代码以支持向量运算。对于Mali T600系列GPU来说，一个向量运算的时间与一个整数加法的时间是一样的。具体代码解读，见下面代码中的注释部分。

[cpp] view plain copy print ? 在CODE上查看代码片 派生到我的代码片
  1. __kernel void hello_world_vector(__global int* restrict inputA,  
  2.                                  __global int* restrict inputB,  
  3.                                  __global int* restrict output)  
  4. {  
  5.     /* 
  6.      * We have reduced the global work size (n) by a factor of 4 compared to the hello_world_opencl sample. 
  7.      * Therefore, i will now be in the range [0, (n / 4) - 1]. 
  8.      */  
  9.     int i = get_global_id(0);  
  10.     /* 
  11.      * Load 4 integers into 'a'. 
  12.      * The offset calculation is implicit from the size of the vector load. 
  13.      * For vloadN(i, p), the address of the first data loaded would be p + i * N. 
  14.      * Load from the data from the address: inputA + i * 4. 
  15.      */  
  16.     int4 a = vload4(i, inputA);  
  17.     /* Do the same for inputB */  
  18.     int4 b = vload4(i, inputB);  
  19.     /* 
  20.      * Do the vector addition. 
  21.      * Store the result at the address: output + i * 4. 
  22.      */  
  23.     vstore4(a + b, i, output);  
  24. }  
　 　由于现在每个内核实例能够实现多个加法运算，所以必须减少内核实例的数量，在宿主机代码中的修改部分如下所示。

[cpp] view plain copy print ? 在CODE上查看代码片 派生到我的代码片
  1. /* 
  2.  * Each instance of our OpenCL kernel now operates on 4 elements of each array so the number of 
  3.  * instances needed is the number of elements in the array divided by 4. 
  4.  */  
  5. size_t globalWorksize[1] = {arraySize / 4};  
  6. /* Enqueue the kernel */  
  7. if (!checkSuccess(clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, globalWorksize, NULL, 0, NULL, &event)))  
  8. {  
  9.     cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);  
  10.     cerr << "Failed enqueuing the kernel. " << __FILE__ << ":"<< __LINE__ << endl;  
  11.     return 1;  
  12. }  
　 　折减系数基于向量的宽度，例如，如果我们在内核中使用int8代替int4，折减系数此时则为8。

　 　

5 运行OpenCL样例


(1). 在SDK根目录的命令行提示符中

[python] view plain copy print ? 在CODE上查看代码片 派生到我的代码片
  1. cd samples\hello_world_vector  
  2. cs-make install  
　 　这样就编译了向量化的OpenCL hello world样例，拷贝了所有运行时需要的文件到SDK根目录下的bin文件夹中。


(2) . 拷贝bin文件夹到目标板中


(3). 在板子上导航到该目录，运行hello world二进制文件

[python] view plain copy print ? 在CODE上查看代码片 派生到我的代码片
  1. chmod 777 hello_world_vector  
  2. ./hello_world_vector 

这篇关于《Mali OpenCL SDK v1.1.0》教程样例之一“Hello World”的文章就介绍到这儿，希望我们推荐的文章对编程师们有所帮助！


http://www.chinasem.cn/article/786399。

相关文章

Spring Security 从入门到进阶系列教程

Spring Security 从入门到进阶系列教程

Spring Security 入门系列 《保护 Web 应用的安全》 《Spring-Security-入门（一）：登录与退出》 《Spring-Security-入门（二）：基于数据库验证》 《Spring-Security-入门（三）：密码加密》 《Spring-Security-入门（四）：自定义-Filter》 《Spring-Security-入门（五）：在 Sprin
阅读更多...
Makefile简明使用教程

Makefile简明使用教程

文章目录 规则makefile文件的基本语法：加在命令前的特殊符号：.PHONY伪目标: Makefilev1 直观写法v2 加上中间过程v3 伪目标v4 变量 make 选项-f-n-C Make 是一种流行的构建工具，常用于将源代码转换成可执行文件或者其他形式的输出文件(如库文件、文档等)。Make 可以自动化地执行编译、链接等一系列操作。 规则 makefile文件
阅读更多...
SWAP作物生长模型安装教程、数据制备、敏感性分析、气候变化影响、R模型敏感性分析与贝叶斯优化、Fortran源代码分析、气候数据降尺度与变化影响分析

SWAP作物生长模型安装教程、数据制备、敏感性分析、气候变化影响、R模型敏感性分析与贝叶斯优化、Fortran源代码分析、气候数据降尺度与变化影响分析

查看原文>>>全流程SWAP农业模型数据制备、敏感性分析及气候变化影响实践技术应用 SWAP模型是由荷兰瓦赫宁根大学开发的先进农作物模型，它综合考虑了土壤-水分-大气以及植被间的相互作用；是一种描述作物生长过程的一种机理性作物生长模型。它不但运用Richard方程，使其能够精确的模拟土壤中水分的运动，而且耦合了WOFOST作物模型使作物的生长描述更为科学。 本文让更多的科研人员和农业工作者
阅读更多...
沁恒CH32在MounRiver Studio上环境配置以及使用详细教程

沁恒CH32在MounRiver Studio上环境配置以及使用详细教程

目录 1.  RISC-V简介 2.  CPU架构现状 3.  MounRiver Studio软件下载 4.  MounRiver Studio软件安装 5.  MounRiver Studio软件介绍 6.  创建工程 7.  编译代码 1.  RISC-V简介         RISC就是精简指令集计算机（Reduced Instruction SetCom
阅读更多...
前端技术(七)——less 教程

前端技术(七)——less 教程

一、less简介 1. less是什么？ less是一种动态样式语言，属于css预处理器的范畴，它扩展了CSS语言，增加了变量、Mixin、函数等特性，使CSS 更易维护和扩展LESS 既可以在 客户端 上运行 ，也可以借助Node.js在服务端运行。 less的中文官网:https://lesscss.cn/ 2. less编译工具 koala 官网 http://koala-app.
阅读更多...
消除安卓SDK更新时的“https://dl-ssl.google.com refused”异常的方法

消除安卓SDK更新时的“https://dl-ssl.google.com refused”异常的方法

消除安卓SDK更新时的“https://dl-ssl.google.com refused”异常的方法   消除安卓SDK更新时的“https://dl-ssl.google.com refused”异常的方法 [转载]原地址：http://blog.csdn.net/x605940745/article/details/17911115 消除SDK更新时的“
阅读更多...
【Shiro】Shiro 的学习教程（三）之 SpringBoot 集成 Shiro

【Shiro】Shiro 的学习教程（三）之 SpringBoot 集成 Shiro

目录 1、环境准备2、引入 Shiro3、实现认证、退出3.1、使用死数据实现3.2、引入数据库，添加注册功能后端代码前端代码 3.3、MD5、Salt 的认证流程 4.、实现授权4.1、基于角色授权4.2、基于资源授权 5、引入缓存5.1、EhCache 实现缓存5.2、集成 Redis 实现 Shiro 缓存 1、环境准备 新建一个 SpringBoot 工程，引入依赖：
阅读更多...
Windows环境利用VS2022编译 libvpx 源码教程

Windows环境利用VS2022编译 libvpx 源码教程

libvpx libvpx 是一个开源的视频编码库，由 WebM 项目开发和维护，专门用于 VP8 和 VP9 视频编码格式的编解码处理。它支持高质量的视频压缩，广泛应用于视频会议、在线教育、视频直播服务等多种场景中。libvpx 的特点包括跨平台兼容性、硬件加速支持以及灵活的接口设计，使其可以轻松集成到各种应用程序中。 libvpx 的安装和配置过程相对简单，用户可以从官方网站下载源代码
阅读更多...
PHP APC缓存函数使用教程

PHP APC缓存函数使用教程

APC，全称是Alternative PHP Cache，官方翻译叫”可选PHP缓存”。它为我们提供了缓存和优化PHP的中间代码的框架。 APC的缓存分两部分:系统缓存和用户数据缓存。(Linux APC扩展安装) 系统缓存 它是指APC把PHP文件源码的编译结果缓存起来，然后在每次调用时先对比时间标记。如果未过期，则使用缓存的中间代码运行。默认缓存 3600s(一小时)。但是这样仍会浪费大量C
阅读更多...
Qt多语种开发教程

Qt多语种开发教程

Qt作为跨平台的开发工具，早已应用到各行各业的软件开发中。 今天讲讲，Qt开发的正序怎么做多语言开发。就是说，你设置中文，就中文显示；设置英语就英文显示，设置繁体就繁体显示，设置发育就显示法语等。 开发环境(其实多语种这块根环境没太大关系)：win10，Qt.5.12.10 一.先用QtCreator创建一个简单的桌面程序 1.工程就随便命名“LanguageTest”，其他默认。 2.在设计师
阅读更多...

Copyright China编程 23002807@qq.com

沪ICP备2023033072号-2