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

1、算法简述


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



2. C/C++实现

　　

[cpp] view plaincopyprint?在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 plaincopyprint?在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 plaincopyprint?在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 plaincopyprint?在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 plaincopyprint?在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 plaincopyprint?在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 plaincopyprint?在CODE上查看代码片派生到我的代码片
  1. for (int i = 0; i < arraySize; i++)  
  2. {  
  3.    inputA[i] = i;  
  4.    inputB[i] = i;  
  5. }  

d. 取消映射缓冲区


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

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


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


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

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

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：http://www.mzph.cn/news/448676.shtml

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈email:809451989@qq.com，一经查实，立即删除！

相关文章

集合去重 （集合元素为引用类型）--- java 8 新特性 --- 根据元素单属性、多属性实现去重

集合去重 （集合元素为引用类型）--- java 8 新特性 --- 根据元素单属性、多属性实现去重

前些天发现了一个巨牛的人工智能学习网站&#xff0c;通俗易懂&#xff0c;风趣幽默&#xff0c;忍不住分享一下给大家。点击跳转到教程。 1. 代码写法&#xff1a; &#xff08;要求 JDK 1.8 或 1.8 以上&#xff09; package gentle.entity;import lombok.Data; /**** auth…
阅读更多...
Django--Forms组件使用

Django--Forms组件使用

Forms组件的使用 在html表单验证中&#xff0c;需要通过各种信息的验证&#xff0c;比如注册界面的姓名、密码、邮箱、电话等的验证&#xff0c;是否符合定义好的规则&#xff0c;不可能每次都要取出对应的字段一一判断&#xff0c;django内置了Forms组件&#xff0c;可以方便的…
阅读更多...
yii2关联表

yii2关联表

asArray()这个方法很好用&#xff0c;返回数组是1版本想要的形式&#xff0c;这种方式有种tp框架的感觉转载于:https://www.cnblogs.com/peipeiyu/p/10974487.html
阅读更多...
详细程序注解学OpenCL一 环境配置和入门程序

详细程序注解学OpenCL一 环境配置和入门程序

本专栏是通过注解程序的方法学习OpenCL&#xff0c;我觉得一个一个地去抠原理也不是办法&#xff0c;干脆直接学习程序&#xff0c;然后把相关原理都直接注解到程序语句当中。 原创地址&#xff1a;http://blog.csdn.net/kenden23/article/details/14101657 一开始要配置好环境…
阅读更多...
解决 IDEA 在 commit 代码时 git 日志乱码 （提交时填写的中文说明乱码）

解决 IDEA 在 commit 代码时 git 日志乱码 （提交时填写的中文说明乱码）

前些天发现了一个巨牛的人工智能学习网站&#xff0c;通俗易懂&#xff0c;风趣幽默&#xff0c;忍不住分享一下给大家。点击跳转到教程。 1.问题描述&#xff1a; idea 开发代码中的 中文正常&#xff0c;但提交到码云时填写的提交日志是乱码。 提交到码云后是这样的&#…
阅读更多...
神奇的pdfkit工具——将字符串保存为pdf文件

神奇的pdfkit工具——将字符串保存为pdf文件

神奇的pdfkit工具——将字符串保存为pdf文件 1、安装工具包 pip install pdfkit 2、上干货 import pdfkitdef create_pdf(str_data, to_file):将字符串生成pdf文件 # &#xff08;需下载wkhtmltox&#xff09;将程序路径传入config对象config pdfkit.configuration(wkhtmltopd…
阅读更多...
OpenCL结构

OpenCL结构

原标题&#xff1a;从零开始学习OpenCL开发&#xff08;一&#xff09;架构 1 异构计算、GPGPU与OpenCL OpenCL是当前一个通用的由很多公司和组织共同发起的多CPU\GPU\其他芯片 异构计算&#xff08;heterogeneous&#xff09;的标准&#xff0c;它是跨平台的。旨在充分利用GP…
阅读更多...
docker-compose 使用小例

docker-compose 使用小例

前些天发现了一个巨牛的人工智能学习网站&#xff0c;通俗易懂&#xff0c;风趣幽默&#xff0c;忍不住分享一下给大家。点击跳转到教程。 只是一个很简单的小例。 1. 原本有的容器 2. docker-compose.yml 写法&#xff1a; gentle 处可以任意写&#xff0c;gentle 是我的项…
阅读更多...
AMD GPU+VS2010的OpenCL配置

AMD GPU+VS2010的OpenCL配置

安装开发环境可以参照DE4-530的OpenCL开发环境搭建&#xff08;最终版&#xff09;&#xff0c;这篇文章的大部分内容转载自&#xff1a;http://www.verydemo.com/demo_c92_i226325.html AMD的Heterogeneous Computing有很多AMD的OpenCL资料&#xff0c;包括各种分析工具&#…
阅读更多...
ABP开发框架前后端开发系列---（9）ABP框架的权限控制管理

ABP开发框架前后端开发系列---（9）ABP框架的权限控制管理

在前面两篇随笔《ABP开发框架前后端开发系列---&#xff08;7&#xff09;系统审计日志和登录日志的管理》和《ABP开发框架前后端开发系列---&#xff08;8&#xff09;ABP框架之Winform界面的开发过程》开始介绍了权限管理的内容&#xff0c;其中只是列出了内部的权限系统的审…
阅读更多...
GIL , 线程池 , 同步 , 异步 , 队列 , 事件

GIL , 线程池 , 同步 , 异步 , 队列 , 事件

一.什么是GIL 官方解释:In CPython, the global interpreter lock, or GIL, is a mutex that prevents multiple native threads from executing Python bytecodes at once. This lock is necessary mainly because CPython’s memory management is not thread-safe. (Howev…
阅读更多...
在Windows下使用OpenCL配置

在Windows下使用OpenCL配置

前言 目前&#xff0c;NVIDIA 和 AMD 的 Windows driver 均有支持OpenCL&#xff08;NVIDIA 的正式版 driver 是从自195.62 版开始&#xff0c;而 AMD则是从9.11 版开始&#xff09;。NVIDIA 的正式版 driver 中包含 OpenCL.dll&#xff0c;因此可以直接使用。AMD 到目前为止…
阅读更多...
CVE-2019-0708 BlueKeep的扫描和打补丁

CVE-2019-0708 BlueKeep的扫描和打补丁

2019独角兽企业重金招聘Python工程师标准>>> 简介 CVE-2019-0708 BlueKeep是一个Windows远程桌面服务的远程代码执行漏洞&#xff0c;其危害程度不亚于CVE-2017-0143 EternalBlue&#xff0c;该漏洞影响了某些旧版本的Windows系统。此漏洞是预身份验证&#xff0c;无…
阅读更多...
《OpenCL异构计算》新版中译本派送中！

《OpenCL异构计算》新版中译本派送中！

《OpenCL异构计算1.2》新鲜出炉&#xff0c;目前市面上仍一书难求&#xff01;我们已向清华出版社订购到第一批新书。关注异构开发社区&#xff0c;积极参与&#xff0c;就有可能免费获取新书&#xff01; 1.如果您异构社区的老朋友&#xff0c;请关注&#xff1a;10.1假期后我…
阅读更多...
查看、关闭当前服务器上启动服务 / 进程

查看、关闭当前服务器上启动服务 / 进程

前些天发现了一个巨牛的人工智能学习网站&#xff0c;通俗易懂&#xff0c;风趣幽默&#xff0c;忍不住分享一下给大家。点击跳转到教程。 我用 java -jar 的方式启动了一个服务&#xff0c;然后要关闭这个服务 / 进程。 1. ps -aux 查看当前进程&#xff0c;整个列表最 后一…
阅读更多...
Booster 系列之——多线程优化

Booster 系列之——多线程优化

项目地址&#xff1a;github.com/didi/booste… 对于开发者来说&#xff0c;线程管理一直是最头疼的问题之一&#xff0c;尤其是业务复杂的 APP&#xff0c;每个业务模块都有着几十甚至上百个线程&#xff0c;而且&#xff0c;作为业务方&#xff0c;都希望本业务的线程优先级最…
阅读更多...
dubbo的底层原理

dubbo的底层原理

前些天发现了一个巨牛的人工智能学习网站&#xff0c;通俗易懂&#xff0c;风趣幽默&#xff0c;忍不住分享一下给大家。点击跳转到教程。 一、Duboo基本概念解释 Dubbo是一种分布式服务框架。 Webservice也是一种服务框架&#xff0c;但是webservice并不是分布式的服务框架&…
阅读更多...
nginx+php+mysql+haproxy+keepalived+NFS，搭建wordpress

nginx+php+mysql+haproxy+keepalived+NFS，搭建wordpress

实现LNMP 实现环境&#xff1a; 服务版本系统CentOS7.6Mysql5.6.34Nginx1.14.2PHP7.1.30HAProxy1.8.20Keepalived1.3.5NFS1.3.0主机IPMysql_master192.168.37.108Mysql_slave192.168.37.105NginxPHP192.168.37.103NginxPHP192.168.37.104HAProxyKeepalived192.168.37.101HAPro…
阅读更多...
Java8函数式编程

Java8函数式编程

最近使用lambda表达式&#xff0c;感觉使用起来非常舒服&#xff0c;箭头函数极大增强了代码的表达能力。于是决心花点时间深入地去研究一下java8的函数式。 一、lambda表达式 先po一个最经典的例子——线程 public static void main(String[] args) {// Java7new Thread(new R…
阅读更多...
IntelliJ IDEA 配置JDK

IntelliJ IDEA 配置JDK

前些天发现了一个巨牛的人工智能学习网站&#xff0c;通俗易懂&#xff0c;风趣幽默&#xff0c;忍不住分享一下给大家。点击跳转到教程。 IDEA配置JDK 1、点击File -->Project Structure&#xff1b; 2、点击左侧标签页SDKs选项&#xff0c;再点击左上角“”&#xff0c;…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023