AMD OpenCL例子阅读笔记系列之AtomicCounters

最新推荐文章于 2023-12-29 10:16:39 发布

清风水岸

最新推荐文章于 2023-12-29 10:16:39 发布

阅读量5k

点赞数 1

分类专栏： OpenCL 文章标签： opencl

本文链接：https://blog.youkuaiyun.com/hanxing0/article/details/24269735

版权

这篇博客记录了作者学习OpenCL过程中对AtomicCounters的理解。文章略过环境配置，重点探讨了OpenCL的扩展项，特别是`#pragma OPENCL EXTENSION`的使用，并对比了全局与非全局变量下`atomic_inc`函数的应用。

摘要生成于 C知道，由 DeepSeek-R1 满血版支持，前往体验 >

学了OpenCL有一段时间了，对于其思想已经有了一些了解。但是对于一些问题仍然难以理解。学习就是要通过阅读代码和自己试验总结经验。所以就有了这个系列的想法。

今天看的时AtomaticCounters例子。这里对环境的配置等就略过了，因为这些可以在网上直接找到。

#pragma OPENCL EXTENSION cl_ext_atomic_counters_32 : enable        

__kernel 
void atomicCounters(
		volatile __global uint *input,
		uint value,
		counter32_t counter)                          
{
	
	size_t globalId = get_global_id(0);
	
	if(value == input[globalId])
		atomic_inc(counter);
		
}                                                                         



/**
 * Counts number of occurrences of value in input array using 
 * global atomics
 */

#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
__kernel 
void globalAtomics(
		volatile __global uint *input,
		uint value,
		__global uint* counter)                         
{                                                                         
	size_t globalId = get_global_id(0);
	
	if(value == input[globalId])
		atomic_inc(&counter[0]);
}

在内核函数中因为要使用到OpenCL中的内置函数，所以在内核中加入了扩展编译选项#pragma OPENCL EXTENSION cl_ext_atomic_counters_32 : enable 以及#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable。先来看下OpenCL扩展，它的格式是：

#pragma OPENCL EXTENSION 扩展项：enable/disable

还有一些比如是否允许化简表达式，格式为：

#pragma OPENCL FP_CONTRACT :ON/OFF/DEFAULT

对于OpenCL的扩展项，其定义的格式为cl_khr_<name>或者cl_<vendor_name>_<name>的开发商形式扩展名。

再看这两个内核实际上区别非常小，主要就是atomic_inc函数的作用在哪里，一个是全局变量，一个是非全局变量，这里看一下在传入内核参数时候的构建方式：

1）atomicCounter:

int
AtomicCounters::runAtomicCounterKernel()
{
    cl_int status = CL_SUCCESS;
    // Set Global and Local work items
    size_t globalWorkItems = length;
    size_t localWorkItems = counterWorkGroupSize;
    // Initialize the counter value
    cl_event writeEvt;
    status = clEnqueueWri

最低0.47元/天解锁文章