OpenCL: 解决图像半透明算法
来源:互联网 发布:alien skin mac 编辑:程序博客网 时间:2024/06/03 16:57
http://www.cocoachina.com/bbs/read.php?tid=33105&keyword=opencl
以下代码就是本人针对之前的OpenCL教程,结合各位能掌握的程度列出的一个示例代码。这份代码由本人亲手炮制。 先贴代码,然后讲解。这里再附上完整的工程。这份代码,当然仍然只能在Snow Leopard下才能运行,因为Leopard上没有OpenCL驱动。
- /*
- * hello.c
- * OpenCL_init
- *
- * Created by Zenny Chen on 9/1/10.
- * Copyright 2010 GreenGames Studio. All rights reserved.
- *
- */
- #include <fcntl.h>
- #include <stdio.h>
- #include <stdlib.h>
- #include <string.h>
- #include <math.h>
- #include <unistd.h>
- #include <sys/types.h>
- #include <sys/stat.h>
- #include <OpenCL/opencl.h>
- ////////////////////////////////////////////////////////////////////////////////
- // Use a static data size for simplicity
- //
- #define IMAGE_X_PIXELS 176
- #define IMAGE_Y_PIXELS 144
- #define IMAGE_SOURCE1_LIMPID 0.5f
- #define IMAGE_SOURCE2_LIMPID 0.5f
- ////////////////////////////////////////////////////////////////////////////////
- // Simple compute kernel which computes the square of an input array
- //
- const char *KernelSource = "\n" \
- "#define IMAGE_Y_PIXELS 144 \n" \
- "#define IMAGE_SOURCE1_LIMPID 0.5f \n" \
- "#define IMAGE_SOURCE2_LIMPID 0.5f \n" \
- " \n" \
- "__kernel void Limpid( \n" \
- " __global float image1[][IMAGE_Y_PIXELS], \n" \
- " __global float image2[][IMAGE_Y_PIXELS], \n" \
- " __global float output[][IMAGE_Y_PIXELS]) \n" \
- "{ \n" \
- " int x = get_global_id(0); \n" \
- " int y = get_global_id(1); \n" \
- " output[x][y] = image1[x][y] * IMAGE_SOURCE1_LIMPID + image2[x][y] * IMAGE_SOURCE2_LIMPID; \n" \
- "} \n" \
- "\n";
- ////////////////////////////////////////////////////////////////////////////////
- int main(int argc, char** argv)
- {
- int err; // error code returned from api calls
- float *image1, *image2; // original data set given to device
- float *results; // results returned from device
- unsigned int correct; // number of correct results returned
- size_t global; // global domain size for our calculation
- size_t local; // local domain size for our calculation
- cl_platform_id platform_id; // added by zenny_chen
- cl_device_id device_id; // compute device id
- cl_context context; // compute context
- cl_command_queue commands; // compute command queue
- cl_program program; // compute program
- cl_kernel kernel; // compute kernel
- cl_mem input1, input2; // device memory used for the input array
- cl_mem output; // device memory used for the output array
- // Initialize the original data buffer and the result buffer
- image1 = (float*)malloc(IMAGE_X_PIXELS * IMAGE_Y_PIXELS * sizeof(*image1));
- image2 = (float*)malloc(IMAGE_X_PIXELS * IMAGE_Y_PIXELS * sizeof(*image2));
- results = (float*)malloc(IMAGE_X_PIXELS * IMAGE_Y_PIXELS * sizeof(*results));
- const unsigned int count = IMAGE_X_PIXELS * IMAGE_Y_PIXELS;
- // Automatically generate 2 images
- for(int i = 0; i < count; i++) {
- image1[i] = rand() / (float)RAND_MAX;
- image2[i] = rand() / (float)RAND_MAX;
- }
- // Create a platform
- err = clGetPlatformIDs(1, &platform_id, NULL);
- if (err != CL_SUCCESS)
- {
- printf("Error: Failed to create a platform!\n");
- return EXIT_FAILURE;
- }
- // Connect to a compute device
- //
- err = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL);
- if (err != CL_SUCCESS)
- {
- printf("Error: Failed to create a device group!\n");
- return EXIT_FAILURE;
- }
- // Create a compute context
- //
- context = clCreateContext((cl_context_properties[]){(cl_context_properties)CL_CONTEXT_PLATFORM, (cl_context_properties)platform_id, 0}, 1, &device_id, NULL, NULL, &err);
- if (!context)
- {
- printf("Error: Failed to create a compute context!\n");
- return EXIT_FAILURE;
- }
- // Create a command commands
- //
- commands = clCreateCommandQueue(context, device_id, 0, &err);
- if (!commands)
- {
- printf("Error: Failed to create a command commands!\n");
- return EXIT_FAILURE;
- }
- // Create the compute program from the source buffer
- //
- program = clCreateProgramWithSource(context, 1, (const char **) & KernelSource, NULL, &err);
- if (!program)
- {
- printf("Error: Failed to create compute program!\n");
- return EXIT_FAILURE;
- }
- // Build the program executable
- //
- err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
- if (err != CL_SUCCESS)
- {
- size_t len;
- char buffer[2048];
- printf("Error: Failed to build program executable!\n");
- clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
- printf("%s\n", buffer);
- exit(1);
- }
- // Create the compute kernel in the program we wish to run
- //
- kernel = clCreateKernel(program, "Limpid", &err);
- if (!kernel || err != CL_SUCCESS)
- {
- printf("Error: Failed to create compute kernel!\n");
- exit(1);
- }
- // Create the input and output arrays in device memory for our calculation
- //
- input1 = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(float) * count, NULL, NULL);
- input2 = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(float) * count, NULL, NULL);
- output = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(float) * count, NULL, NULL);
- if (!input1 || !input2 || !output)
- {
- printf("Error: Failed to allocate device memory!\n");
- exit(1);
- }
- // Write our data set into the input array in device memory
- //
- err = clEnqueueWriteBuffer(commands, input1, CL_TRUE, 0, sizeof(float) * count, image1, 0, NULL, NULL);
- err |= clEnqueueWriteBuffer(commands, input2, CL_TRUE, 0, sizeof(float) * count, image2, 0, NULL, NULL);
- if (err != CL_SUCCESS)
- {
- printf("Error: Failed to write to source array!\n");
- exit(1);
- }
- // Set the arguments to our compute kernel
- //
- err = 0;
- err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &input1);
- err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &input2);
- err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &output);
- if (err != CL_SUCCESS)
- {
- printf("Error: Failed to set kernel arguments! %d\n", err);
- exit(1);
- }
- // Get the maximum work group size for executing the kernel on the device
- //
- err = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local), &local, NULL);
- if (err != CL_SUCCESS)
- {
- printf("Error: Failed to retrieve kernel work group info! %d\n", err);
- exit(1);
- }
- else
- printf("The number of work items in a work group is: %lu\r\n", local);
- // Execute the kernel over the entire range of our 1d input data set
- // using the maximum number of work group items for this device
- //
- global = count;
- err = clEnqueueNDRangeKernel(commands, kernel, 2, NULL, (size_t[]){IMAGE_X_PIXELS, IMAGE_Y_PIXELS}, (size_t[]){22, 12}, 0, NULL, NULL);
- if (err)
- {
- printf("Error: Failed to execute kernel!\n");
- return EXIT_FAILURE;
- }
- // Wait for the command commands to get serviced before reading back results
- //
- clFinish(commands);
- // Read back the results from the device to verify the output
- //
- err = clEnqueueReadBuffer(commands, output, CL_TRUE, 0, sizeof(float) * count, results, 0, NULL, NULL );
- if (err != CL_SUCCESS)
- {
- printf("Error: Failed to read output array! %d\n", err);
- exit(1);
- }
- // Validate our results
- //
- correct = 0;
- for(int i = 0; i < count; i++)
- {
- if(results[i] == image1[i] * IMAGE_SOURCE1_LIMPID + image2[i] * IMAGE_SOURCE2_LIMPID)
- correct++;
- }
- // Print a brief summary detailing the results
- //
- printf("Computed '%d/%d' correct values!\n", correct, count);
- // Shutdown and cleanup
- //
- clReleaseMemObject(input1);
- clReleaseMemObject(input2);
- clReleaseMemObject(output);
- clReleaseProgram(program);
- clReleaseKernel(kernel);
- clReleaseCommandQueue(commands);
- clReleaseContext(context);
- free(image1);
- free(image2);
- free(results);
- return 0;
- }
附件: OpenCL_init.zip (18 K) 下载次数:279
这里架构上与第二讲的Apple提供的示例代码差不多。但是,这边的内核函数要比之前的复杂一些。 这里输入有两个缓存数组,并且用了二维工作维度。我们首先看一下Limpid内核计算函数。 OpenCL中的C语言是ISO/IEC9899:1999标准的子集,因此兼容大多数C99的语法特性。我们在定义一个二维数组的时候必须要指定列数,函数可以忽略。这是C语言的一个特性。具体原因,可以在我的CSDN博客中找到答案。 这里的Limpid内核函数有两个输入数组,分别用来存放两个图像数据。我们分别对两个图像相应的每对像素做插值混合计算,将结果像素数据送给输出缓存。这里的像素是用一个单精度浮点来表示的;当然,我们也可以看作为是由4个单精度浮点组成。 这里我们通过调用get_global_id()来获得当前工作项第i维的id,对于二维工作组,每个二维ID(id_x, id_y)用于唯一标识一个工作项线程所对应的存储单元。对于二维工作组而言,第0维表示数据的列索引;而第1维则表示数据的行索引。如果我们将二维数组用一维索引来表示的话就是id_x + id_y * x_length;x_length表示一行有多少列。 我们分别获得当前行、列索引就能确定唯一的存储单元。我们可以对此进行操作。
然后我们再看第213行,对clEnqueueNDRangeKernel函数的调用。这里我们指定工作维度是2。然后,全局工作项的个数就分别是图像长度和高度;而一个工作组的大小这里是需要好好讲一下的。 由于我们之前通过clGetKernelWorkGroupInfo函数获得的工作组最大大小是512,而如果我们只用一个工作组的话就需要176 * 144个工作项,这个乘积已经大大超过512了。因此,我们可以对176 * 144进行划分,使得每个工作组的实际大小小于512。我们这里对一个工作组的行取了22,列取了12。那么对于二维工作组而言,工作组的个数就被划分为(8, 12)个。如果将整个计算资源看作为一个大的容器的话,那么它就有8 * 12个工作组;每个工作组有22 * 12个工作项。 如果各位有CUDA编程的一些概念的话,那么上面所说的工作组其实对应于Block;而上面所说的容器,其实就是一个Grid。这里再要提醒一下的是全局工作项的每个维度的个数都必须能够被所划分后的一个工作组的每个维度的工作项个数整除。
我们在主机端做数据组织的话可以不用考虑维度,都用一维的也没问题。包括数据创建以及读写。
下面为该程序加入了Profiling,并加入了1024 * 768的尺寸。在1024*768的情况下,工作组的X向可以含有32个工作项;Y向可以含有16个工作项。这样,X * Y正好等于512,能充分发挥GPU的效能。而此时,工作组集合就被划分成了(32, 48)。
附件: OpenCL_Transparent.zip (26 K) 下载次数:42
下面提供更好的一个工程。这个工程里,OpenCL的内核做成文件资源,并且通过XCode能够显示出语法高亮,这样,内核代码就更容易阅读和修改,呵呵。 这里由于也是基于控制台的,因此在添加file copy的时候应该鼠标右键Copy Files,然后选择Get Info,将属性设置为 Resource,并在path里面空白。这样,在构造整个工程时,XCode会将你的kernel.cl作为资源文件拷贝到目标可执行文件的目录中去。
附件: OpenCL_Basic.zip (20 K) 下载次数:42
- OpenCL: 解决图像半透明算法
- OpenCL: 解决图像半透明算法
- 基于OpenCL的图像积分图算法实现
- 基于OpenCL的图像积分图算法改进
- 基于FPGA的Alpha半透明图像叠加算法实现
- 半透明算法
- 半透明算法
- 半透明算法:
- OpenCL嵌入式图像处理
- OpenCL编程:图像卷积
- 图像半透明的实现
- OpenCL实现MapReduce算法
- 遮罩层内容半透明解决
- 机器学习决策树算法解决图像识别
- opencl+opencv实现sobel算法
- jQuery 解决 ie6 半透明 插件
- opencl:C++实现双线性插值图像缩放
- 图像旋转在opencL 实现方法
- jbpm4.4下hibernate映射文件中表跟类的对应关系
- 网络爬虫
- HDU 2095(map)
- 范数规则化
- 仿path菜单的实现原理,例ArcMenu
- OpenCL: 解决图像半透明算法
- cut、tee、split、xargs、bc命令
- Ubuntu安装配置NFS文件系统
- cocos2d 竖屏
- spring mvc 与struts2区别
- 网络金融让银行出一身冷汗
- 一切为了老婆
- android开发之activity的跳转关闭问题
- C# 命令行编译器详解