Apple's OpenCL——再谈Local Memory (转)

http://blog.csdn.net/zenny_chen/article/details/6159746

在OpenCL中，用__local（或local）修饰的变量会被存放在一个计算单元（Compute Unit）的共享存储器区域中。对于nVidia的GPU，一个CU可以被映射为物理上的一块SM（Stream Multiprocessor）；而对于AMD-ATi的GPU可以被映射为物理上的一块SIMD。不管是SM也好，SIMD也罢，它们都有一个在本计算单元中被所有线程（OpenCL中称为Work Item）所共享的共享存储器。因此，在一个计算单元内，可以通过local shared memory来同步此计算单元内的所有工作项。

这里必须注意的是在计算单元之间的线程的通信只能通过全局存储器进行，因为每个计算单元之间是没有共享存储器的，呵呵。

 

下面我将证明Apple的OpenCL实现中，如果有两个Work Group（一个Work Group的处理交给一个计算单元执行），那么这两个Work Group正好能分别被映射到一个计算单元内。我用的是Mac Mini，GPU为GeForce 9400M，所有仅有两个SM，呵呵。

下面先给出kernel代码：

[cpp] view plaincopyprint?
  1. __kernel void solve_sum( 
  2.                     __global volatile unsigned buffer[512], 
  3.                     __global unsigned dest[512] 
  4.                     ) 
  5. { 
  6.     __local volatileint flag = 0; 
  7.      
  8.     size_t gid = get_global_id(0); 
  9.      
  10.     if(0 <= gid && gid < 32) 
  11.     { 
  12.         while(flag != 1); 
  13.         flag = 0; 
  14.          
  15.         buffer[gid] = 0x1UL; 
  16.         //write_mem_fence(CLK_GLOBAL_MEM_FENCE); 
  17.     } 
  18.     elseif(32 <= gid && gid < 64) 
  19.     { 
  20.         flag = 1; 
  21.          
  22.         while(flag != 0); 
  23.         unsigned ret = buffer[31 + 32 - gid]; 
  24.         dest[gid - 32] = ret; 
  25.     } 
  26. } 

 

上面这个内核程序的配置为：分为两个工作组；每组32个工作项。这样，两个工作组能进不同的SM。各位在执行这段代码时会发生死循环。然后等2到3秒后程序会自动退出，这点不用担心，呵呵。原因就是两个SM的共享变量flag是各有各的一份。假定，线程0到线程31进SM0，那么SM0的所有线程共享这个flag变量；而线程32到线程63进SM1，那么SM1的flag将被SM1的所有线程共享。而如果企图把这个（其实是两个）共享变量用于两个SM之间的通信，显然是无法成功的，呵呵。尽管代码上只写了一个flag，但实际上却有两个副本。

 

下面提供主机端代码：

[cpp] view plaincopyprint?
  1. #import <Foundation/Foundation.h> 
  2. #include <OpenCL/opencl.h> 
  3. static unsigned __attribute__((aligned(16))) buffer[512] = { 0 };    // original data set given to device 
  4. static unsigned __attribute__((aligned(16))) dest[512] = { 0 }; 
  5. int opencl_execution(void) 
  6. { 
  7.     int err;                            // error code returned from api calls 
  8.      
  9.     size_t local;                       // local domain size for our calculation 
  10.      
  11.     cl_platform_id  platform_id;        // added by zenny_chen 
  12.     cl_device_id device_id;             // compute device id  
  13.     cl_context context;                 // compute context 
  14.     cl_command_queue commands;          // compute command queue 
  15.     cl_program program;                 // compute program 
  16.     cl_kernel kernel;                   // compute kernel 
  17.      
  18.     cl_mem memOrg, memDst;                      // device memory used for the input array 
  19.      
  20.      
  21.     // Create a platform 
  22.     err = clGetPlatformIDs(1, &platform_id, NULL); 
  23.     if (err != CL_SUCCESS) 
  24.     { 
  25.         printf("Error: Failed to create a platform!/n"); 
  26.         return EXIT_FAILURE; 
  27.     } 
  28.      
  29.     // Connect to a compute device 
  30.     // 
  31.     err = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL); 
  32.     if (err != CL_SUCCESS) 
  33.     { 
  34.         printf("Error: Failed to create a device group!/n"); 
  35.         return EXIT_FAILURE; 
  36.     } 
  37.      
  38.     // Create a compute context  
  39.     // 
  40.     context = clCreateContext((cl_context_properties[]){(cl_context_properties)CL_CONTEXT_PLATFORM, (cl_context_properties)platform_id, 0}, 1, &device_id, NULL, NULL, &err); 
  41.     if (!context) 
  42.     { 
  43.         printf("Error: Failed to create a compute context!/n"); 
  44.         return EXIT_FAILURE; 
  45.     } 
  46.      
  47.     // Create a command commands 
  48.     // 
  49.     commands = clCreateCommandQueue(context, device_id, 0, &err); 
  50.     if (!commands) 
  51.     { 
  52.         printf("Error: Failed to create a command commands!/n"); 
  53.         return EXIT_FAILURE; 
  54.     } 
  55.      
  56.     // Fetch kernel source 
  57.     NSString *filepath = [[NSBundle mainBundle] pathForResource:@"kernel" ofType:@"cl"]; 
  58.     if(filepath == NULL) 
  59.     { 
  60.         puts("Source not found!"); 
  61.         return EXIT_FAILURE; 
  62.     } 
  63.      
  64.     constchar *KernelSource = (constchar*)[[NSString stringWithContentsOfFile:filepath encoding:NSUTF8StringEncoding error:nil] UTF8String]; 
  65.      
  66.     // Create the compute program from the source buffer 
  67.     // 
  68.     program = clCreateProgramWithSource(context, 1, (constchar **) & KernelSource, NULL, &err); 
  69.     if (!program) 
  70.     { 
  71.         printf("Error: Failed to create compute program!/n"); 
  72.         return EXIT_FAILURE; 
  73.     } 
  74.      
  75.     // Build the program executable 
  76.     // 
  77.     err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL); 
  78.     if (err != CL_SUCCESS) 
  79.     { 
  80.         size_t len; 
  81.         char buffer[2048]; 
  82.          
  83.         printf("Error: Failed to build program executable!/n"); 
  84.         clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len); 
  85.         printf("%s/n", buffer); 
  86.         exit(1); 
  87.     } 
  88.      
  89.     // Create the compute kernel in the program we wish to run 
  90.     // 
  91.     kernel = clCreateKernel(program, "solve_sum", &err); 
  92.     if (!kernel || err != CL_SUCCESS) 
  93.     { 
  94.         printf("Error: Failed to create compute kernel!/n"); 
  95.         exit(1); 
  96.     } 
  97.      
  98.     // Create the input and output arrays in device memory for our calculation 
  99.     // 
  100.     memOrg = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int) * 512, NULL, NULL); 
  101.     memDst = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(int) * 512, NULL, NULL); 
  102.      
  103.     if (memOrg == NULL || memDst == NULL) 
  104.     { 
  105.         printf("Error: Failed to allocate device memory!/n"); 
  106.         exit(1); 
  107.     }     
  108.      
  109.     // Write our data set into the input array in device memory  
  110.     // 
  111.     err = clEnqueueWriteBuffer(commands, memOrg, CL_TRUE, 0, sizeof(int) * 512, buffer, 0, NULL, NULL); 
  112.     if (err != CL_SUCCESS) 
  113.     { 
  114.         printf("Error: Failed to write to source array!/n"); 
  115.         exit(1); 
  116.     } 
  117.      
  118.     // Set the arguments to our compute kernel 
  119.     // 
  120.     err = 0; 
  121.     err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &memOrg); 
  122.     err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &memDst); 
  123.     if (err != CL_SUCCESS) 
  124.     { 
  125.         printf("Error: Failed to set kernel arguments! %d/n", err); 
  126.         exit(1); 
  127.     } 
  128.      
  129.     // Get the maximum work group size for executing the kernel on the device 
  130.     // 
  131.     err = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local), &local, NULL); 
  132.     if (err != CL_SUCCESS) 
  133.     { 
  134.         printf("Error: Failed to retrieve kernel work group info! %d/n", err); 
  135.         exit(1); 
  136.     } 
  137.     else 
  138.         printf("The number of work items in a work group is: %lu/r/n", local); 
  139.      
  140.     // Execute the kernel over the entire range of our 1d input data set 
  141.     // using the maximum number of work group items for this device 
  142.     // 
  143.      
  144.     err = clEnqueueNDRangeKernel(commands, kernel, 1, NULL, (size_t[]){ 64 }, (size_t[]){ 32 }, 0, NULL, NULL); 
  145.     if (err) 
  146.     { 
  147.         printf("Error: Failed to execute kernel!/n"); 
  148.         return EXIT_FAILURE; 
  149.     } 
  150.      
  151.     // Wait for the command commands to get serviced before reading back results 
  152.     // 
  153.     clFinish(commands); 
  154.      
  155.     // Read back the results from the device to verify the output 
  156.     // 
  157.     err = clEnqueueReadBuffer(commands, memDst, CL_TRUE, 0, sizeof(int) * 512, dest, 0, NULL, NULL );   
  158.     if (err != CL_SUCCESS) 
  159.     { 
  160.         printf("Error: Failed to read output array! %d/n", err); 
  161.         exit(1); 
  162.     } 
  163.      
  164.     // Validate our results 
  165.     // 
  166.     printf("The result is: 0x%.8X/n", dest[0]); 
  167.      
  168.      
  169.     // Shutdown and cleanup 
  170.     // 
  171.     clReleaseMemObject(memOrg); 
  172.     clReleaseMemObject(memDst); 
  173.     clReleaseProgram(program); 
  174.     clReleaseKernel(kernel); 
  175.     clReleaseCommandQueue(commands); 
  176.     clReleaseContext(context); 
  177.      
  178.     return 0; 
  179. } 
  180. int main (int argc, constchar * argv[]) { 
  181.     NSAutoreleasePool * pool = [[NSAutoreleasePool alloc] init]; 
  182.     // insert code here... 
  183.     opencl_execution(); 
  184.     [pool drain]; 
  185.     return 0; 
  186. } 

 

见主机端代码第144行：

[cpp] view plaincopyprint?
  1. err = clEnqueueNDRangeKernel(commands, kernel, 1, NULL, (size_t[]){ 64 }, (size_t[]){ 32 }, 0, NULL, NULL);  

 

这里，我们设定全局工作项个数为64，每个工作组有32个线程，那么这样一来就自然地被划分为两个工作组。如果我们把32改为64，这么一来就变为一个工作组，这样，在一个SM中通过一个共享变量进行通信完全OK，程序就能正常终止。

另外，如果想保持原来的2个Work Group，那么必须通过全局变量进行通信：

[cpp] view plaincopyprint?
  1. __kernel void solve_sum( 
  2.                     __global volatile unsigned buffer[512], 
  3.                     __global unsigned dest[512] 
  4.                     ) 
  5. { 
  6.     __local volatileint flag = 0; 
  7.      
  8.     size_t gid = get_global_id(0); 
  9.      
  10.     if(0 <= gid && gid < 32) 
  11.     { 
  12.         while(buffer[256] != 1); 
  13.         buffer[256] = 0; 
  14.          
  15.         buffer[gid] = 0x1UL; 
  16.         //write_mem_fence(CLK_GLOBAL_MEM_FENCE); 
  17.     } 
  18.     elseif(32 <= gid && gid < 64) 
  19.     { 
  20.         buffer[256] = 1; 
  21.          
  22.         while(buffer[256] != 0); 
  23.         unsigned ret = buffer[31 + 32 - gid]; 
  24.         dest[gid - 32] = ret; 
  25.     } 
  26. } 

 

这里还要注意一点。用于通信的变量都必须加上volatile，否则，OpenCL内核编译器会把对全局变量的第二次访问全都优化为直接从寄存器取数据，从而外部对此变量的改变在当前线程内将无法看见。

posted @ 2012-03-26 17:21  董雨  阅读(426)  评论(0)  编辑  收藏  举报