opencl:cl::make_kernel的进化

已于 2024-01-02 10:55:03 修改
阅读量2.5k 收藏 4
点赞数 1
分类专栏： c/c++/c++11 opencl 文章标签： opencl kernel C++11 变长模板 模板函数
于 2016-03-10 12:18:19 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/10km/article/details/50832741
版权

我之前的一篇博客《opencl:C++ 利用cl::make_kernel简化kernel执行代码》详细说明了如何使用OpenCL C++接口(cl.hpp)提供cl::make_kernel算子来简化kernel执行代码。

/* 缩放图像(双线性插值) */
gray_matrix_cl gray_matrix_cl::zoom(size_t dst_width, size_t dst_height, const facecl_context& context)const {
	gray_matrix_cl dst_matrix(dst_width, dst_height);
	auto command_queue = global_facecl_context.getCommandQueue();// 获取cl::CommandQueue
	this->upload(command_queue);//向OpenCL设备中上传原始图像数据
	cl_float widthNormalizationFactor = 1.0f / dst_width;
	cl_float heightNormalizationFactor = 1.0f / dst_height;
	//构造cl::make_kernel对象执行kernel
	cl::make_kernel<cl::Image2D,cl::Image2D,cl_float,cl_float>
		(context.getKernel(KERNEL_NAME(image_scaling)))// 获取已经编译好的cl::Kernel
		(cl::EnqueueArgs(command_queue,cl::NDRange( dst_width, dst_height )),
		cl_img,dst_matrix.cl_img,
		widthNormalizationFactor,
		heightNormalizationFactor);
	command_queue.finish(); // 等待kernel执行结束
	dst_matrix.download(command_queue);从OpenCL设备中下载结果数据
	return std::move(dst_matrix);
}

这是上一篇博客中最后简化的代码。与原来原始代码相比，这种调用方式将所有设置kernel参数的调用(setArg)都被cl::make_kernel算子(fuctor)封装，调用者不需要知道细节。只需要执行cl::make_kernel的operator(),在()中按kernel定义的参数顺序将kernel需要的参数填在括号中，cl::make_kernel算子会自动为kernel设置参数并将kernel压入command_queue执行。

Ok,前一篇博客的内容回顾完毕。
那么还能不能进一步改进，让kernel执行更简单化？
再看看上面的代码，在用opencl的kernel执行一个图像的缩放之前，先要

this->upload(command_queue);//向OpenCL设备中上传原始图像数据

在kernel执行结束之后，

dst_matrix.download(command_queue);从OpenCL设备中下载结果数据

在你写完第一个kernel程序后，再写另外一个kernel的时候,你会发现几乎所有的kernel调用都要有上面两个动作，概括起来就是

  1. 在执行kernel之前，如果kernel参数中有指针类型或imag类型的参数，需要将参数在主机端对应的cl::Memory类型(其子类包括cl::Image,cl::Buffer)的数据上传(upload)到设备
  2. 在执行kernel结束后，可能需要将kernel处理之后的输出数据(同样是cl::Memory类型)下载(download)到主机。

这些都是重复和类似的代码，我们只要把这两个动作抽象出来（memory_cl类），就可以有办法将这两个动作也封装起来。

关于如何实现memory_cl类，将要本文后面讲到，现在假定我们已经有memory_cl类实现对所有cl::Memory对象的download和upload统一管理

make_kernel进化之run_kernel

于是利用C++11的变长模板特性，我们可以写出下面的run_kernel模板函数

template<typename IN_CL_TYPE // kernel参数中的输入数据类型(cl::Buffer,cl::Image)
		,typename OUT_CL_TYPE// kernel参数中的输出数据类型(cl::Buffer,cl::Image) 
		,typename... Args    // kernel参数中其他标量数据类型，变长模板，允许多个参数
		>
void run_kernel(const cl::EnqueueArgs &queue_args // 队列参数
		,const cl::Kernel &kernel//kernel对象
		,bool download           //kernel执行结束后是否将结果数据下载到本地？  
		,const memory_cl<IN_CL_TYPE> &in // 输入数据对象，memory_cl为自已写的opencl内存管理类
		,memory_cl<OUT_CL_TYPE>&out// 输出数据对象，memory_cl为自已写的opencl内存管理类
		,Args&&... args //其他kernel参数
		){
	// 根据数据状态标记判断是否需要上传数据到设备，如果数据已经在设备中就不需要upload
	in.upload_if_need(queue_args.queue_);
	// 执行kernel
	cl::make_kernel<IN_CL_TYPE, OUT_CL_TYPE, Args...>k(kernel);//创建cl::make_kernel对象
	k(queue_args,in.cl_mem_obj,out.cl_mem_obj, std::forward<Args>(args)...);//执行kernel
	// 根据download标记决定是否执行	memory_cl的download函数将kernel输出数据下载到主机。
	if(download)
		out.download(queue_args.queue_);
}

借助这个run_kernel模板函数，前面实现图像缩放的gray_matrix_cl::zoom函数就可以改写如下：

template<typename T>T get_align(T v,uint8_t a){return (T)((v+(T)((1<<a)-1))>>a);}
/* 缩放图像(双线性插值) */
gray_matrix_cl gray_matrix_cl::zoom(size_t dst_width, size_t dst_height, const facecl_context& context,bool download)const {
	gray_matrix_cl dst_matrix(dst_width, dst_height);
	auto command_queue = global_facecl_context.getCommandQueue();
	cl_float widthNormalizationFactor = 1.0f / dst_width;
	cl_float heightNormalizationFactor = 1.0f / dst_height;
	run_kernel(
			cl::EnqueueArgs(command_queue,{ 1, get_align(dst_height,4) })//队列参数对象
			,context.getKernel(KERNEL_NAME(image_scaling)) // 要执行的kernel对象
			,true //自动下载结果数据
			,*this //输入图像
			,dst_matrix // 输出图像
			,widthNormalizationFactor
			,heightNormalizationFactor
		);
	command_queue.finish(); // 等待kernel执行结束
	return std::move(dst_matrix);
}

哈哈，这样以来代码又简化了，大功告成！

run_kernel进化

但是好像当我准备将这个run_kernel，用于执行第二个kernel函数时，问题来了。
我们看上面这个run_kernel函数，它对kernel函数的参数类型和顺序是有要求的：

  1. 第一个参数必须是输入的数据对象
  2. 第二个参数必须是输出数据对象
  3. 其他标量数据对象必须位于第三位以后

所以，它的使用是有限制的，我的第二个kernel函数，只有一个数据对象参数，它即是输入又是输出，它就不太方便用这个函数，(当然还是可以用，将这参数重复填入两次)
当kernel函数有超一个输入数据对象或输出数据对象，就没可能用这个模板函数。。。

能不能改进run_kernel函数，使它允许接收超过一个输入/出数据对象参数，并且不用限定kernel的参数顺序呢？
yes,we can
run_kernel要经历再一次的进化！

下面是改进后的run_kernel模板函数

template<typename... Args>
inline void run_kernel_new(const cl::EnqueueArgs &queue_args// 队列参数对象
		, const cl::Kernel &kernel // kernel对象
		, bool download // kernel执行结束后是否下载结果数据
		, Args&&... args //  kernel参数表
		){
	// 根据需要上传所有cl::Memory对象的数据到设备
	upload_args_if_need<1>(queue_args.queue_,std::forward<Args>(args)...);
	typename make_make_kernel<Args...>::type k(kernel);
	k(queue_args,std::forward<Args>(args)...); // 执行kernel
	// 根据download标记需要下载所有cl::Memory输出对象的数据到主机
	download_args<1>(queue_args.queue_,download,std::forward<Args>(args)...);
}

额，粗看起来与前一版本的run_kernel,貌似差不多，
但还是它真的是进化了

进化之一

只是参数中不再有in,out参数，也就是说，参数表中可以不用关心in/out参数的顺序以及个数了。

		,const memory_cl<IN_CL_TYPE> &in
		,memory_cl<OUT_CL_TYPE>&out

进化之二

与前一版本的run_kernel相比，原来第一行的in.upload_if_need(queue_args.queue_);换成了upload_args_if_need<1>(queue_args.queue_,std::forward<Args>(args)...);最后一行的out.download(queue_args.queue_);换成了download_args<1>(queue_args.queue_,download,std::forward<Args>(args)...);

等等， 这upload_args_if_need和download_args是个模板函数啊，
嗯，在这里用了递归模板函数，循环检查args 参数表中的参数类型，如果是memory_cl类就执行memory_cl中的upload_if_need函数，
download_args也是差不多，如果是memory_cl类就根据download标记执行memory_cl中的download函数
upload_args_if_need和download_args模板函数的实现如下：

/* 模板函数，检查T是否为memory_cl的子类 */
template<typename T>
struct is_kind_of_memory_cl{
	template <typename CL_TYPE>
		static CL_TYPE  check(memory_cl<CL_TYPE>);
	static void check(...);
	using cl_type=decltype(check(std::declval<T>()));
	enum{value=!std::is_same<cl_type,void>::value};
};
/*
 * upload_arg(x)_if_need和download_arg(x)系列模板函数循环对run_kernel中的所有变长参数类型进行识别,
 * 对于memory_cl类型的参数，根据需要在kernel执行前上传数据到设备，
 * 并在kernel执行后根据需要下载输出数据到主机
 * 模板中的N参数，用于调试时知道哪个参数出错
 *
 * */
// 参数ARG为非memory_cl类型时直接返回，啥也不做
template<int N,typename ARG>
typename std::enable_if<!is_kind_of_memory_cl<ARG>::value>::type
inline upload_arg_if_need(const cl::CommandQueue &command_queue,const ARG & arg){}
// 参数ARG是memory_cl类型,时根据需要上传数据
template<int N,typename ARG>
typename std::enable_if<is_kind_of_memory_cl<ARG>::value>::type
inline upload_arg_if_need(const cl::CommandQueue &command_queue,const ARG & arg){
	const cl::Memory&m=arg.cl_mem_obj;
	auto mem_context=m.getInfo<CL_MEM_CONTEXT>();
	auto queue_context=command_queue.getInfo<CL_QUEUE_CONTEXT>();
	// 检查memory_cl中内存对象的context与command_queue是否一致，不一致则抛出异常
	if(mem_context()!=queue_context()){
		std::stringstream stream;
		stream<<":the arg No:"<<N;// 动态参数编号
		throw std::invalid_argument(std::string(SOURCE_AT).append(stream.str()).append(":mem_context()!=queue_context()"));
	}
	try{
		arg.upload_if_need(command_queue);//上传数据到设备
	}catch(cl::Error&e){
		std::stringstream stream;
		stream<<"the arg No:"<<N;// 动态参数编号
		throw face_cl_exception(SOURCE_AT,e,stream.str());
	}catch(face_exception&e){
		std::stringstream stream;
		stream<<"the arg No:"<<N<<e.what();// 动态参数编号
		throw face_cl_exception(SOURCE_AT,stream.str());
	}catch(std::exception&e){
		std::stringstream stream;
		stream<<"the arg No:"<<N;// 动态参数编号
		throw face_cl_exception(SOURCE_AT,e,stream.str());
	}catch(...){
		std::stringstream stream;
		stream<<"the arg No:"<<N<<":unknow exception";// 动态参数编号
		throw face_cl_exception(SOURCE_AT,stream.str());
	}
}
// 特例:参数表为空，递归终止
template<int N>
inline void upload_args_if_need(const cl::CommandQueue &command_queue){
}
/* 递归处理Args中的每一个参数
 * 如果是memory_cl类型的对象，则上传数据到设备
 * */
template<int N,typename ARG1,typename... Args>
inline void upload_args_if_need(const cl::CommandQueue &command_queue,ARG1 && arg1,Args&&... args){
	upload_arg_if_need<N>	(command_queue,std::forward<ARG1>(arg1));//处理第一个参数
	upload_args_if_need<N+1>	(command_queue,std::forward<Args>(args)...);//递归处理其他参数
}
// 参数ARG为非memory_cl类型时，为空函数，啥也不做直接返回
template<int N,typename ARG>
typename std::enable_if<!is_kind_of_memory_cl<ARG>::value>::type
inline download_arg(const cl::CommandQueue &command_queue,bool download, const ARG & arg){}
// 参数ARG是memory_cl类型,时根据需要下载数据到主机
template<int N,typename ARG>
typename std::enable_if<is_kind_of_memory_cl<ARG>::value>::type
inline download_arg(const cl::CommandQueue &command_queue,bool download, const ARG & arg){
	if(download){
		try{
			const cl::Memory &m=arg.cl_mem_obj;
			auto flags=m.getInfo<CL_MEM_FLAGS>();
			// 根据CL_MEM_FLAGS判断是否为输出数据对象，以决定是否需要下载数据
			if(flags&(CL_MEM_WRITE_ONLY|CL_MEM_READ_WRITE)){
				const_cast<ARG&>(arg).download(command_queue);//下载数据到设备
			}
		}catch(cl::Error&e){
			std::stringstream stream;
			stream<<"the arg No:"<<N;// 动态参数编号
			throw face_cl_exception(SOURCE_AT,e,stream.str());
		}catch(face_exception&e){
			std::stringstream stream;
			stream<<"the arg No:"<<N<<e.what();// 动态参数编号
			throw face_cl_exception(SOURCE_AT,stream.str());
		}catch(std::exception&e){
			std::stringstream stream;
			stream<<"the arg No:"<<N;// 动态参数编号
			throw face_cl_exception(SOURCE_AT,e,stream.str());
		}catch(...){
			std::stringstream stream;
			stream<<"the arg No:"<<N<<":unknow exception";// 动态参数编号
			throw face_cl_exception(SOURCE_AT,stream.str());
		}
	}
}
// 特例:参数表为空，递归终止
template<int N>
inline void download_args(const cl::CommandQueue &command_queue,bool download){}
/* 递归处理Args中的每一个参数
 * 如果是memory_cl类型的对象，则根据download参数的指示下载数据到主机
 * */
template<int N,typename ARG1,typename... Args>
inline void download_args(const cl::CommandQueue &command_queue,bool download, ARG1 && arg1,Args&&... args){
	download_arg<N>(command_queue,download,std::forward<ARG1>(arg1));//处理第一个参数
	download_args<N+1>(command_queue,download,std::forward<Args>(args)...);//递归处理其他参数
}

进化之三

原来是直接实例化cl::make_kernel类对象的

	cl::make_kernel<IN_CL_TYPE, OUT_CL_TYPE, Args...>k(kernel);

而新版本则改成了

typename make_make_kernel<Args...>::type k(kernel);

这里make_make_kernel也是一个模板函数，用来实例化cl::make_kernel类，为什么要这么做呢？

因为传递给run_kernel的参数中所有OpenCL内存对象(cl::Buffer,cl::Image)都被我自定义的memeory_cl类封装起来了，而cl::make_kernel在执行的时候，参数类型却是需要原始的OpenCL内存对象(cl::Buffer,cl::Image)，所以实例化cl::make_kernel时必须将memeory_cl类型转为对应的OpenCL内存对象类型。
make_make_kernel模板函数就是实现这个功能的，下面是make_make_kernel的代码实现


/* 模板函数返回make_kernel执行里需要的类
 * 对于普通的类，就是类本身
 * 对于memory_cl的子类，返回memory_cl::cl_cpp_type
 *  */
template<typename ARG
		,typename ARG_TYPE=typename std::decay<ARG>::type
		,typename MEM_CL= is_kind_of_memory_cl<ARG>
		,typename K_TYPE=typename std::conditional<MEM_CL::value,typename MEM_CL::cl_type,ARG>::type
		>
struct kernel_type {
	using type= K_TYPE;
};

/*
 * 模板函数
 * 根据模板参数，创建cl::make_kernel类
 * 创建cl::make_kernel类时所有的模板参数都会调用 kernel_type模板函数,
 * 以获取实例化cl::make_kernel时真正需要的类型
*/
template <
   typename T0,   typename T1 = cl::detail::NullType,   typename T2 = cl::detail::NullType,
   typename T3 = cl::detail::NullType,   typename T4 = cl::detail::NullType,
   typename T5 = cl::detail::NullType,   typename T6 = cl::detail::NullType,
   typename T7 = cl::detail::NullType,   typename T8 = cl::detail::NullType,
   typename T9 = cl::detail::NullType,   typename T10 = cl::detail::NullType,
   typename T11 = cl::detail::NullType,   typename T12 = cl::detail::NullType,
   typename T13 = cl::detail::NullType,   typename T14 = cl::detail::NullType,
   typename T15 = cl::detail::NullType,   typename T16 = cl::detail::NullType,
   typename T17 = cl::detail::NullType,   typename T18 = cl::detail::NullType,
   typename T19 = cl::detail::NullType,   typename T20 = cl::detail::NullType,
   typename T21 = cl::detail::NullType,   typename T22 = cl::detail::NullType,
   typename T23 = cl::detail::NullType,   typename T24 = cl::detail::NullType,
   typename T25 = cl::detail::NullType,   typename T26 = cl::detail::NullType,
   typename T27 = cl::detail::NullType,   typename T28 = cl::detail::NullType,
   typename T29 = cl::detail::NullType,   typename T30 = cl::detail::NullType,
   typename T31 = cl::detail::NullType
>
struct make_make_kernel{
	using type=cl::make_kernel<
			typename kernel_type<T0>::type,		typename kernel_type<T1>::type,
			typename kernel_type<T2>::type,		typename kernel_type<T3>::type,
			typename kernel_type<T4>::type,		typename kernel_type<T5>::type,
			typename kernel_type<T6>::type,		typename kernel_type<T7>::type,
			typename kernel_type<T8>::type,		typename kernel_type<T9>::type,
			typename kernel_type<T10>::type,	typename kernel_type<T11>::type,
			typename kernel_type<T12>::type,	typename kernel_type<T13>::type,
			typename kernel_type<T14>::type,	typename kernel_type<T15>::type,
			typename kernel_type<T16>::type,	typename kernel_type<T17>::type,
			typename kernel_type<T18>::type,	typename kernel_type<T19>::type,
			typename kernel_type<T20>::type,	typename kernel_type<T21>::type,
			typename kernel_type<T22>::type,	typename kernel_type<T23>::type,
			typename kernel_type<T24>::type,	typename kernel_type<T25>::type,
			typename kernel_type<T26>::type,	typename kernel_type<T27>::type,
			typename kernel_type<T28>::type,	typename kernel_type<T29>::type,
			typename kernel_type<T30>::type,	typename kernel_type<T31>::type
			>;
};

总结

进化后的run_kernel使用起来了方便多了，对kernel参数个数和顺序不再有限制，同时自动实现OpenCL内存对象数据的上传和下载。
只是代码貌似增加了好多好多，实现增加的代码主要是模板函数，都只是在编译期起作用，并不会增加多少运行时代码。
它带来的好处是当你的项目中有很多不同的kernel函数要执行时，使用这种设计方式可以大大减少撰写重复或相似的代码，同时增加代码的稳定性。

神奇的memory_cl

前面一直不断被提起的用来封装OpenCL内存对象的memory_cl是个什么神奇的东东？呵呵，其实并不复杂，就是抽象的基类而已，下面是这个类的主要实现代码和函数声明。前面代码所涉及到的所有函数都在这里有声明。

/*
 * OpenCL内存抽象模型定义
 * memory_cl为抽象接口,所有OpenCL内存对象(cl::Buffer,cl::Image等等)都被封装在该对象内部
 * 主要提供主机与设备之间的交换功能
 * 项目中涉及的其他涉及OpenCL内存对象的类都是此类的衍生类
 * matrix_cl 继承自memory_cl，是抽象矩阵类
 * integral_matrix继承自matrix_cl，积分图对象类
 * gray_matrix_cl继承自matrix_cl,灰度图像类
 * */
template<typename CL_TYPE,
		typename ENABLE=typename std::enable_if<std::is_base_of<cl::Memory,CL_TYPE>::value>::type>
class memory_cl{
public:
	using cl_cpp_type=CL_TYPE;
private:
	mutable bool	on_device=false;	// 数据是否已经在设备上标志
public:
	cl_cpp_type cl_mem_obj; // OpenCL 内存对象
	/* 如果数据没有上传到设备(on_device=false)，则向OpenCL设备中上传原始矩阵数据,
	 * 上传成功则将on_device置为true
	 * */
	void upload_if_need(const cl::CommandQueue& command_queue=Null_Queue)const{
		if(!on_device){
			upload(command_queue);
		}
	}

	/* 虚函数，从OpenCL设备中下载结果数据, 将on_device标志置为true */
	virtual void download(const cl::CommandQueue& command_queue=Null_Queue){
		throw face_exception(SOURCE_AT,"sub class must implement the funtion "
				"by calling download_force(const cl::CommandQueue& command_queue,std::vector<E> &out)");
	}
	/* 虚函数，向OpenCL设备中上传原始矩阵数据, 将on_device标志置为true */
	virtual void upload(const cl::CommandQueue& command_queue=Null_Queue)const{
		throw face_exception(SOURCE_AT,
				"sub class must implement the funtion "
				"by calling upload_force(const cl::CommandQueue& command_queue,std::vector<E> &in) ");
	}
	/* upload_force上传cl::Memory对象到设备,上传成功则将on_device置为true
	 * 因为项目中只涉及到使用cl::Buffer和cl::Image2D所以,在此做只分别对cl::Buffer和cl::Image写了相关的代码，
	 * download_force也是一样
	 */
	template<typename E, typename _CL_TYPE = CL_TYPE>
	typename std::enable_if<std::is_base_of<cl::Buffer,_CL_TYPE>::value>::type
	upload_force(const std::vector<E> &in,const cl::CommandQueue& command_queue=Null_Queue) const;
	template<typename E,typename _CL_TYPE=CL_TYPE>
	typename std::enable_if<std::is_base_of<cl::Image2D,_CL_TYPE>::value>::type
	upload_force(const std::vector<E> &in,const cl::CommandQueue& command_queue=Null_Queue) const;
	/* 从cl_mem_obj对象中下载数据到out,下载成功则将on_device置为true */
	template<typename E, typename _CL_TYPE = CL_TYPE>
	typename std::enable_if<std::is_base_of<cl::Buffer,_CL_TYPE>::value>::type
	download_force(std::vector<E> &out, const cl::CommandQueue& command_queue=Null_Queue) const;
	template<typename E, typename _CL_TYPE = CL_TYPE>
	typename std::enable_if<std::is_base_of<cl::Image2D,_CL_TYPE>::value>::type
	download_force(std::vector<E> &out,size_t row_pitch=0,const cl::CommandQueue& command_queue=Null_Queue) const;
	//相关的构造函数/
	memory_cl(const CL_TYPE& cl_mem_obj,bool on_device):cl_mem_obj(cl_mem_obj),on_device(on_device){};
	memory_cl(const memory_cl&)=default;
	memory_cl(memory_cl&&)=default;
	memory_cl()=default;
	memory_cl& operator=(const memory_cl&)=default;
	memory_cl& operator=(memory_cl&&rv){
		this->cl_mem_obj=std::move(rv.cl_mem_obj);
		this->on_device=rv.on_device;
		return *this;
	};
	/* operator type()操作符，返回OpenCL内存对象 */
	operator const cl_cpp_type& ()const{	return this->cl_mem_obj;	}
	operator cl_cpp_type&(){return this->cl_mem_obj;}
	virtual ~memory_cl()=default;
};
10km
关注
  • 1
    点赞
  • 踩
  • 4
    收藏
    觉得还不错? 一键收藏
  • 打赏
    打赏
  • 1
    评论
专栏目录
OpenCL 图像处理函数、图像放缩和插值
INFINALGEORGE的博客
04-03 4150
1.OpenCL提供了大量可以在内核中运行的图像处理函数，它们大致可以分为以下三类： （1）Read functiongs--返回给定坐标上的颜色取值； （2）write functiongs-- 设定给定坐标上的颜色取值； （3）Information functions-- 提供关于图像对象的信息，例如图像的维度以及像素属性； 图像读取函数是从图像对象中读取向量，他们各自的参数基本一样。唯一...
OpenCL应用与双线性插值
SE_and_CS
03-28 3165
一．  CPU实现和GPU实现 本实验基于的测试图像是8000*1000大小的8位灰度钢板图像，缩小倍数为0.5。 CPU上算法实现的核心代码 void func(BYTE*src, int width, int height, int pitchSrc, BYTE *dst, int t_width, int t_height, int pitchDes) { float f
OpenCL专题02：Makefile案例解析
吴智深的博客
06-21 369
Makefile案例解析 Makefile文件基本结构 递归依赖 常用字符串处理函数
端侧GPU基于opencl实现reduce算子
LuchangLi 的专栏
08-03 977
端侧GPU基于opencl实现reduce算子
CLION openCL C/C++环境配置(两个坑)
zmh657628320的博客
11-05 967
CLION 配置 OpenCL环境 基于书籍opencl in action代码
opencl:C++实现双线性插值图像缩放
10km的专栏
02-27 8333
用OpenCL实现图像缩放代码是比较简单的，因为OpenCL本身就支持双线性插值 下面是kernel代码(从Mali OpenCL SDK 抄来的：/samples/image_scaling/assets/image_scaling.cl) 非常简单只有4行// 定义采样器 // CLK_NORMALIZED_COORDS_TRUE指定使用归一化坐标 // CLK_ADDRESS_CLAMP
香菜：为OpenCL:trade_mark:1.2设备构建NVIDIA:registered:CUDA:trade_mark:代码
02-06
为OpenCL:trade_mark:1.2设备构建以NVIDIA:registered:CUDA:trade_mark:代码编写的应用程序。 概念 将应用程序留在NVIDIA:registered:CUDA:trade_mark:中 编译成OpenCL 1.2 在任何OpenCL 1.2 GPU上运行 如何使用 ...
ros_opencl:ROS + OpenCL =
05-11
标题"ros_opencl:ROS + OpenCL ="暗示了这个项目是将ROS与OpenCL结合，以利用OpenCL的并行计算能力优化ROS中的计算密集型任务。这可能是为了提高机器人的实时性能，尤其是在处理视觉数据、SLAM（Simultaneous ...
GlfwCL: Tiny GLFW OpenCL GL & stb_image:GlfwCL：使用GLFW和stb_image的小型可扩展OpenCL / GL项目-开源
05-12
**GlfwCL** 是一个基于 **GLFW** 和 **OpenCL** 的小型开源项目，它同时结合了 **OpenGL** 和 **stb_image** 库，为开发者提供了一个简洁且易于扩展的框架，用于构建高性能的图形和计算应用程序。GLFW 是一个流行的...
OpenCL-ICD-Loader:OpenCL ICD加载程序项目
04-29
OpenCL ICD加载器此存储库... 您需要为支持OpenCL ICD扩展cl_khr_icd OpenCL设备获取并安装OpenCL实施，才能使用OpenCL ICD加载程序运行应用程序。 OpenCL扩展规范中描述了OpenCL可安装客户端驱动程序扩展（ cl_khr_icd
kernel_slicer:kernel_slicer
03-31
kernel_slicer：CPU至GPU（Vulkan）转换器 图1.由于许多现有的编程模式，我们的翻译有很多头。 该项目基于clang（图1）。 我们衷心感谢clang前端开发人员的出色结构和项目文档！ 项目概况 kernel_slicer是原型...
OpenCL学习笔记（一）开发环境搭建（win10+vs2019）
最新发布
aggs1990的专栏
12-08 1940
异构编程开发，在高性能编程中有重要的，笔者本次只简单介绍下，如何搭建简单的开发环境，可以供有需要的小伙伴们开发测试使用。
【高性能计算】opencl安装及相关概念
种一棵树最好的时间是十年前，其次是现在
08-24 1780
安装特定硬件厂商提供的OpenCL SDK：如果您的系统不提供默认的OpenCL运行时和驱动程序，或者您希望使用特定硬件厂商提供的最新版本运行时和驱动程序，那么可以安装相应厂商提供的OpenCL SDK。虽然特定硬件厂商提供的OpenCL SDK可能包含额外的工具和优化，但OpenCL的核心标准和函数接口是跨平台的，并且可以在没有特定SDK的情况下进行开发和编译。总而言之，异构计算利用不同类型的计算资源的特点和优势，通过合理分配和协同使用这些资源，实现更高的计算性能、能源效率和灵活性。
OpenCL 学习step by step (5) 使用二维NDRange workgroup
weixin_34037173的博客
09-07 328
      在本教程中，我们使用二维NDRange来设置workgroup，这样在opencl中，workitme的组织形式是二维的，Kernel中 的代码也要做相应的改变，我们先看一下clEnqueueNDRangeKernel函数的变化。首先我们指定了workgroup size为localx*localy,通常这个值为64的倍数，但最好不要超过256。 //执行kernel,Range用2...
OpenCL的NDRange中globalSize，globalSize参数理解及调试笔记
固本培元的专栏
03-13 6789
先上源码： const size_t globalSize[2] = {6,4} ; const size_t localSize[2] = {3,2} ; const size_t offset[2] = {3,5} ; size_t dim = 2 ; errNum = clEnqueueNDRangeKernel(commandQueue, kern
使用二维NDRange workgroup
weixin_34090562的博客
01-08 193
作为初学者一直，经过多次的上网搜索你一定会看到迈克老狼的向量加法的示例，不知道你是否和我一样，刚开始并不是很准确的知道他的add.cl写的代码的意思，源码如下： #pragma OPENCL EXTENSION cl_amd_printf : enable __kernel void vecadd(__global const float* a, __global const float*...
cmake总结
浪子荆的博客
01-29 336
学习了makefile的规则之后，感觉makefile的规则还是有些复杂的，要建立一个比较复杂的项目，makefile还是有一定的工作量的。但是现在已经有了很多工具来将makefile给封装起来，减少工作量。现在要介绍的一个就是cmake，很多开源的项目也是使用cmake的，还有一个autotool工具，实现类似的功能。下面的文章就是介绍cmake的使用的。 1.http://www.hahac...
CMake实例
Salted的博客
09-28 186
makefile: 这里是将它作为运行cmake的脚本进行使用 cmake : -mkdir ./build cd ./build && cmake .. -DCMAKE_BUILD_TYPE=Debug CMakeList.txt cmake_minimum_required(VERSION 3.10) #当前cmake版本 set(CMAKE_CXX_STANDARD 14)# -std=c++17 project(first_proj) set(CMAKE_EXPORT_.
写一段C++代码，把opencv读入的图像cv::Mat A转化为cl_mem格式送入opencl的命令队列
03-09
我可以回答这个问题。以下是将opencv读入的图像cv::Mat A转化为cl_mem格式送入opencl的命令队列的C代码示例： ``` // 假设已经创建了OpenCL上下文和命令队列 // 假设cv::Mat A是一个3通道的图像 // 创建OpenCL内存对象 cl_mem cl_A = clCreateBuffer(context, CL_MEM_READ_ONLY, A.total() * A.elemSize(), NULL, &err); // 将cv::Mat A的数据复制到OpenCL内存对象中 err = clEnqueueWriteBuffer(command_queue, cl_A, CL_TRUE, , A.total() * A.elemSize(), A.data, , NULL, NULL); // 将OpenCL内存对象添加到OpenCL内核参数中 err = clSetKernelArg(kernel, , sizeof(cl_mem), (void *)&cl_A); // 执行OpenCL内核 err = clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, global_work_size, local_work_size, , NULL, NULL); ```

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论 1
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

10km

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付：¥1
获取中
扫码支付

您的余额不足，请更换扫码支付或充值

打赏作者

实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值