#include <flow_graph_opencl_node.h>

Inheritance diagram for tbb::flow::interface11::opencl_factory< DeviceFilter >:

Collaboration diagram for tbb::flow::interface11::opencl_factory< DeviceFilter >:

Classes
struct	finalize_fn

struct	finalize_fn_leaf

class	kernel

Public Types
template<typename T >
using	async_msg_type = opencl_async_msg< T, opencl_factory< DeviceFilter >>

typedef opencl_device	device_type

typedef kernel	kernel_type

typedef opencl_range	range_type

Public Member Functions
	opencl_factory ()

	~opencl_factory ()

bool	init (const opencl_device_list &device_list)

template<typename... Args>
void	send_kernel (opencl_device device, const kernel_type &kernel, const range_type &work_size, Args &...args)

template<typename T , typename... Rest>
void	send_data (opencl_device device, T &t, Rest &...args)

void	send_data (opencl_device)

template<typename FinalizeFn , typename... Args>
void	finalize (opencl_device device, FinalizeFn fn, Args &...args)

const opencl_device_list &	devices ()

Private Member Functions
template<typename Factory >
void	enqueue_map_buffer (opencl_device device, opencl_buffer_impl< Factory > &buffer, opencl_async_msg< void *, Factory > &dmsg)

template<typename Factory >
void	enqueue_unmap_buffer (opencl_device device, opencl_memory< Factory > &memory, opencl_async_msg< void *, Factory > &dmsg)

template<size_t NUM_ARGS, typename T >
void	process_one_arg (const kernel_type &kernel, std::array< cl_event, NUM_ARGS > &, int &, int &place, const T &t)

template<size_t NUM_ARGS, typename T , typename F >
void	process_one_arg (const kernel_type &kernel, std::array< cl_event, NUM_ARGS > &events, int &num_events, int &place, const opencl_async_msg< T, F > &msg)

template<size_t NUM_ARGS, typename T , typename... Rest>
void	process_arg_list (const kernel_type &kernel, std::array< cl_event, NUM_ARGS > &events, int &num_events, int &place, const T &t, const Rest &...args)

template<size_t NUM_ARGS>
void	process_arg_list (const kernel_type &, std::array< cl_event, NUM_ARGS > &, int &, int &)

template<typename T >
void	update_one_arg (cl_event, T &)

template<typename T , typename F >
void	update_one_arg (cl_event e, opencl_async_msg< T, F > &msg)

template<typename T , typename... Rest>
void	update_arg_list (cl_event e, T &t, Rest &...args)

void	update_arg_list (cl_event)

cl_event	send_kernel_impl (opencl_device device, const cl_kernel &kernel, const range_type &work_size, cl_uint num_events, cl_event *event_list)

template<typename T >
bool	get_event_from_one_arg (cl_event &, const T &)

template<typename T , typename F >
bool	get_event_from_one_arg (cl_event &e, const opencl_async_msg< T, F > &msg)

template<typename T , typename... Rest>
bool	get_event_from_args (cl_event &e, const T &t, const Rest &...args)

bool	get_event_from_args (cl_event &)

bool	is_same_context (opencl_device::device_id_type d1, opencl_device::device_id_type d2)

	opencl_factory (const opencl_factory &)

opencl_factory &	operator= (const opencl_factory &)

cl_context	context ()

void	init_once ()

Static Private Member Functions
static void CL_CALLBACK	finalize_callback (cl_event, cl_int event_command_exec_status, void *data)

Private Attributes
std::once_flag	my_once_flag

opencl_device_list	my_devices

cl_context	my_cl_context

tbb::spin_mutex	my_devices_mutex

Friends
template<typename Factory >
class	opencl_program

template<typename Factory >
class	opencl_buffer_impl

template<typename Factory >
class	opencl_memory

Detailed Description

template<typename DeviceFilter>
class tbb::flow::interface11::opencl_factory< DeviceFilter >

Definition at line 59 of file flow_graph_opencl_node.h.

Member Typedef Documentation

template<typename DeviceFilter>

template<typename T >

using tbb::flow::interface11::opencl_factory< DeviceFilter >::async_msg_type = opencl_async_msg<T, opencl_factory<DeviceFilter>>

Definition at line 826 of file flow_graph_opencl_node.h.

template<typename DeviceFilter>

typedef opencl_device tbb::flow::interface11::opencl_factory< DeviceFilter >::device_type

Definition at line 827 of file flow_graph_opencl_node.h.

template<typename DeviceFilter>

typedef kernel tbb::flow::interface11::opencl_factory< DeviceFilter >::kernel_type

Definition at line 869 of file flow_graph_opencl_node.h.

template<typename DeviceFilter>

typedef opencl_range tbb::flow::interface11::opencl_factory< DeviceFilter >::range_type

Definition at line 873 of file flow_graph_opencl_node.h.

Constructor & Destructor Documentation

template<typename DeviceFilter>

tbb::flow::interface11::opencl_factory< DeviceFilter >::opencl_factory ( )

inline

Definition at line 875 of file flow_graph_opencl_node.h.

875 {}

template<typename DeviceFilter>

tbb::flow::interface11::opencl_factory< DeviceFilter >::~opencl_factory ( )

inline

Definition at line 876 of file flow_graph_opencl_node.h.

                       {
         if ( my_devices.size() ) {
             for ( auto d = my_devices.begin(); d != my_devices.end(); ++d ) {
                 enforce_cl_retcode( clReleaseCommandQueue( (*d).my_cl_command_queue ), "Failed to release a command queue" );
             }
             enforce_cl_retcode( clReleaseContext( my_cl_context ), "Failed to release a context" );
         }
     }

template<typename DeviceFilter>

tbb::flow::interface11::opencl_factory< DeviceFilter >::opencl_factory ( const opencl_factory< DeviceFilter > & )

private

Member Function Documentation

template<typename DeviceFilter>

cl_context tbb::flow::interface11::opencl_factory< DeviceFilter >::context ( )

inlineprivate

Definition at line 1092 of file flow_graph_opencl_node.h.

                          {
         std::call_once( my_once_flag, &opencl_factory::init_once, this );
         return my_cl_context;
     }

template<typename DeviceFilter>

const opencl_device_list& tbb::flow::interface11::opencl_factory< DeviceFilter >::devices ( )

inline

Definition at line 1075 of file flow_graph_opencl_node.h.

                                         {
         std::call_once( my_once_flag, &opencl_factory::init_once, this );
         return my_devices;
     }

template<typename DeviceFilter>

template<typename Factory >

void tbb::flow::interface11::opencl_factory< DeviceFilter >::enqueue_map_buffer	(	opencl_device	device,
		opencl_buffer_impl< Factory > &	buffer,
		opencl_async_msg< void *, Factory > &	dmsg
	)

inlineprivate

Definition at line 897 of file flow_graph_opencl_node.h.

                                                                                                                                  {
         cl_event const* e1 = dmsg.get_event();
         cl_event e2;
         cl_int err;
         void *ptr = clEnqueueMapBuffer( device.my_cl_command_queue, buffer.get_cl_mem(), false, CL_MAP_READ | CL_MAP_WRITE, 0, buffer.size(),
             e1 == NULL ? 0 : 1, e1, &e2, &err );
         enforce_cl_retcode( err, "Failed to map a buffer" );
         dmsg.data( false ) = ptr;
         dmsg.set_event( e2 );
         enforce_cl_retcode( clReleaseEvent( e2 ), "Failed to release an event" );
     }

template<typename DeviceFilter>

template<typename Factory >

void tbb::flow::interface11::opencl_factory< DeviceFilter >::enqueue_unmap_buffer	(	opencl_device	device,
		opencl_memory< Factory > &	memory,
		opencl_async_msg< void *, Factory > &	dmsg
	)

inlineprivate

Definition at line 911 of file flow_graph_opencl_node.h.

                                                                                                                               {
         cl_event const* e1 = dmsg.get_event();
         cl_event e2;
         enforce_cl_retcode(
             clEnqueueUnmapMemObject( device.my_cl_command_queue, memory.get_cl_mem(), memory.get_host_ptr(), e1 == NULL ? 0 : 1, e1, &e2 ),
            "Failed to unmap a buffer" );
         dmsg.set_event( e2 );
         enforce_cl_retcode( clReleaseEvent( e2 ), "Failed to release an event" );
     }

template<typename DeviceFilter>

template<typename FinalizeFn , typename... Args>

void tbb::flow::interface11::opencl_factory< DeviceFilter >::finalize	(	opencl_device	device,
		FinalizeFn	fn,
		Args &...	args
	)

inline

Definition at line 1064 of file flow_graph_opencl_node.h.

                                                                         {
         cl_event e;
 
         if ( get_event_from_args( e, args... ) ) {
             enforce_cl_retcode( clSetEventCallback( e, CL_COMPLETE, finalize_callback,
                 new finalize_fn_leaf<FinalizeFn>(fn) ), "Failed to set a callback" );
         }
 
         enforce_cl_retcode( clFlush( device.my_cl_command_queue ), "Failed to flush an OpenCL command queue" );
     }

template<typename DeviceFilter>

static void CL_CALLBACK tbb::flow::interface11::opencl_factory< DeviceFilter >::finalize_callback	(	cl_event	,
		cl_int	event_command_exec_status,
		void *	data
	)

inlinestaticprivate

Definition at line 1051 of file flow_graph_opencl_node.h.

                                                                                                       {
         tbb::internal::suppress_unused_warning(event_command_exec_status);
         __TBB_ASSERT(event_command_exec_status == CL_COMPLETE, NULL);
 
         finalize_fn * const fn_ptr = static_cast<finalize_fn*>(data);
         __TBB_ASSERT(fn_ptr != NULL, "Invalid finalize function pointer");
         (*fn_ptr)();
 
         // Function pointer was created by 'new' & this callback must be called once only
         delete fn_ptr;
     }

template<typename DeviceFilter>

template<typename T , typename... Rest>

bool tbb::flow::interface11::opencl_factory< DeviceFilter >::get_event_from_args	(	cl_event &	e,
		const T &	t,
		const Rest &...	args
	)

inlineprivate

Definition at line 1026 of file flow_graph_opencl_node.h.

                                                                              {
         if ( get_event_from_one_arg( e, t ) ) {
             return true;
         }
 
         return get_event_from_args( e, args... );
     }

template<typename DeviceFilter>

bool tbb::flow::interface11::opencl_factory< DeviceFilter >::get_event_from_args ( cl_event & )

inlineprivate

Definition at line 1034 of file flow_graph_opencl_node.h.

                                           {
         return false;
     }

template<typename DeviceFilter>

template<typename T >

bool tbb::flow::interface11::opencl_factory< DeviceFilter >::get_event_from_one_arg	(	cl_event &	,
		const T &
	)

inlineprivate

Definition at line 1009 of file flow_graph_opencl_node.h.

                                                        {
         return false;
     }

template<typename DeviceFilter>

template<typename T , typename F >

bool tbb::flow::interface11::opencl_factory< DeviceFilter >::get_event_from_one_arg	(	cl_event &	e,
		const opencl_async_msg< T, F > &	msg
	)

inlineprivate

Definition at line 1014 of file flow_graph_opencl_node.h.

                                                                                  {
         cl_event const *e_ptr = msg.get_event();
 
         if ( e_ptr != NULL ) {
             e = *e_ptr;
             return true;
         }
 
         return false;
     }

template<typename DeviceFilter>

bool tbb::flow::interface11::opencl_factory< DeviceFilter >::init ( const opencl_device_list & device_list )

inline

Definition at line 885 of file flow_graph_opencl_node.h.

                                                        {
         tbb::spin_mutex::scoped_lock lock( my_devices_mutex );
         if ( !my_devices.size() ) {
             my_devices = device_list;
             return true;
         }
         return false;
     }

template<typename DeviceFilter>

void tbb::flow::interface11::opencl_factory< DeviceFilter >::init_once ( )

inlineprivate

Definition at line 1097 of file flow_graph_opencl_node.h.

Referenced by tbb::flow::interface11::opencl_factory< default_device_filter >::context(), and tbb::flow::interface11::opencl_factory< default_device_filter >::devices().

                      {
         {
             tbb::spin_mutex::scoped_lock lock(my_devices_mutex);
             if (!my_devices.size())
                 my_devices = DeviceFilter()( opencl_info::available_devices() );
         }
 
         enforce_cl_retcode(my_devices.size() ? CL_SUCCESS : CL_INVALID_DEVICE, "No devices in the device list");
         cl_platform_id platform_id = my_devices.begin()->platform_id();
         for (opencl_device_list::iterator it = ++my_devices.begin(); it != my_devices.end(); ++it)
             enforce_cl_retcode(it->platform_id() == platform_id ? CL_SUCCESS : CL_INVALID_PLATFORM, "All devices should be in the same platform");
 
         std::vector<cl_device_id> cl_device_ids;
         for (auto d = my_devices.begin(); d != my_devices.end(); ++d) {
             cl_device_ids.push_back((*d).my_cl_device_id);
         }
 
         cl_context_properties context_properties[3] = { CL_CONTEXT_PLATFORM, (cl_context_properties)platform_id, (cl_context_properties)NULL };
         cl_int err;
         cl_context ctx = clCreateContext(context_properties,
             (cl_uint)cl_device_ids.size(),
             cl_device_ids.data(),
             NULL, NULL, &err);
         enforce_cl_retcode(err, "Failed to create context");
         my_cl_context = ctx;
 
         size_t device_counter = 0;
         for (auto d = my_devices.begin(); d != my_devices.end(); d++) {
             (*d).my_device_id = device_counter++;
             cl_int err2;
             cl_command_queue cq;
 #if CL_VERSION_2_0
             if ((*d).major_version() >= 2) {
                 if ((*d).out_of_order_exec_mode_on_host_present()) {
                     cl_queue_properties props[] = { CL_QUEUE_PROPERTIES, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, 0 };
                     cq = clCreateCommandQueueWithProperties(ctx, (*d).my_cl_device_id, props, &err2);
                 } else {
                     cl_queue_properties props[] = { 0 };
                     cq = clCreateCommandQueueWithProperties(ctx, (*d).my_cl_device_id, props, &err2);
                 }
             } else
 #endif
             {
                 cl_command_queue_properties props = (*d).out_of_order_exec_mode_on_host_present() ? CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE : 0;
                 // Suppress "declared deprecated" warning for the next line.
 #if __TBB_GCC_WARNING_SUPPRESSION_PRESENT
 #pragma GCC diagnostic push
 #pragma GCC diagnostic ignored "-Wdeprecated-declarations"
 #endif
 #if _MSC_VER || __INTEL_COMPILER
 #pragma warning( push )
 #if __INTEL_COMPILER
 #pragma warning (disable: 1478)
 #else
 #pragma warning (disable: 4996)
 #endif
 #endif
                 cq = clCreateCommandQueue(ctx, (*d).my_cl_device_id, props, &err2);
 #if _MSC_VER || __INTEL_COMPILER
 #pragma warning( pop )
 #endif
 #if __TBB_GCC_WARNING_SUPPRESSION_PRESENT
 #pragma GCC diagnostic pop
 #endif
             }
             enforce_cl_retcode(err2, "Failed to create command queue");
             (*d).my_cl_command_queue = cq;
         }
     }

Here is the caller graph for this function:

template<typename DeviceFilter>

bool tbb::flow::interface11::opencl_factory< DeviceFilter >::is_same_context	(	opencl_device::device_id_type	d1,
		opencl_device::device_id_type	d2
	)

inlineprivate

Definition at line 1081 of file flow_graph_opencl_node.h.

                                                                                            {
         __TBB_ASSERT( d1 != opencl_device::unknown && d2 != opencl_device::unknown, NULL );
         // Currently, factory supports only one context so if the both devices are not host it means the are in the same context.
         if ( d1 != opencl_device::host && d2 != opencl_device::host )
             return true;
         return d1 == d2;
     }

template<typename DeviceFilter>

opencl_factory& tbb::flow::interface11::opencl_factory< DeviceFilter >::operator= ( const opencl_factory< DeviceFilter > & )

private

template<typename DeviceFilter>

template<size_t NUM_ARGS, typename T , typename... Rest>

void tbb::flow::interface11::opencl_factory< DeviceFilter >::process_arg_list	(	const kernel_type &	kernel,
		std::array< cl_event, NUM_ARGS > &	events,
		int &	num_events,
		int &	place,
		const T &	t,
		const Rest &...	args
	)

inlineprivate

Definition at line 941 of file flow_graph_opencl_node.h.

                                                                                                                                                            {
         process_one_arg( kernel, events, num_events, place, t );
         process_arg_list( kernel, events, num_events, place, args... );
     }

template<typename DeviceFilter>

template<size_t NUM_ARGS>

void tbb::flow::interface11::opencl_factory< DeviceFilter >::process_arg_list	(	const kernel_type &	,
		std::array< cl_event, NUM_ARGS > &	,
		int &	,
		int &
	)

inlineprivate

Definition at line 947 of file flow_graph_opencl_node.h.

947 {}

template<typename DeviceFilter>

template<size_t NUM_ARGS, typename T >

void tbb::flow::interface11::opencl_factory< DeviceFilter >::process_one_arg	(	const kernel_type &	kernel,
		std::array< cl_event, NUM_ARGS > &	,
		int &	,
		int &	place,
		const T &	t
	)

inlineprivate

Definition at line 923 of file flow_graph_opencl_node.h.

                                                                                                                    {
         auto p = get_native_object(t);
         enforce_cl_retcode( clSetKernelArg(kernel.my_cl_kernel, place++, sizeof(p), &p), "Failed to set a kernel argument" );
     }

template<typename DeviceFilter>

template<size_t NUM_ARGS, typename T , typename F >

void tbb::flow::interface11::opencl_factory< DeviceFilter >::process_one_arg	(	const kernel_type &	kernel,
		std::array< cl_event, NUM_ARGS > &	events,
		int &	num_events,
		int &	place,
		const opencl_async_msg< T, F > &	msg
	)

inlineprivate

Definition at line 929 of file flow_graph_opencl_node.h.

                                                                                                                                                             {
         __TBB_ASSERT((static_cast<typename std::array<cl_event, NUM_ARGS>::size_type>(num_events) < events.size()), NULL);
 
         const cl_event * const e = msg.get_event();
         if (e != NULL) {
             events[num_events++] = *e;
         }
 
         process_one_arg( kernel, events, num_events, place, msg.data(false) );
     }

template<typename DeviceFilter>

template<typename T , typename... Rest>

void tbb::flow::interface11::opencl_factory< DeviceFilter >::send_data	(	opencl_device	device,
		T &	t,
		Rest &...	args
	)

inline

Definition at line 983 of file flow_graph_opencl_node.h.

                                                               {
         send_if_memory_object( device, t );
         send_data( device, args... );
     }

template<typename DeviceFilter>

void tbb::flow::interface11::opencl_factory< DeviceFilter >::send_data ( opencl_device )

inline

Definition at line 988 of file flow_graph_opencl_node.h.

988 {}

template<typename DeviceFilter>

template<typename... Args>

void tbb::flow::interface11::opencl_factory< DeviceFilter >::send_kernel	(	opencl_device	device,
		const kernel_type &	kernel,
		const range_type &	work_size,
		Args &...	args
	)

inline

Definition at line 967 of file flow_graph_opencl_node.h.

                                                                                                                     {
         std::array<cl_event, sizeof...(Args)> events;
         int num_events = 0;
         int place = 0;
         process_arg_list( kernel, events, num_events, place, args... );
 
         const cl_event e = send_kernel_impl( device, kernel.my_cl_kernel, work_size, num_events, events.data() );
 
         update_arg_list(e, args...);
 
         // Release our own reference to cl_event
         enforce_cl_retcode( clReleaseEvent(e), "Failed to release an event" );
     }

template<typename DeviceFilter>

cl_event tbb::flow::interface11::opencl_factory< DeviceFilter >::send_kernel_impl	(	opencl_device	device,
		const cl_kernel &	kernel,
		const range_type &	work_size,
		cl_uint	num_events,
		cl_event *	event_list
	)

inlineprivate

Definition at line 992 of file flow_graph_opencl_node.h.

                                                                                 {
         const typename range_type::nd_range_type g_offset = { { 0, 0, 0 } };
         const typename range_type::nd_range_type& g_size = work_size.global_range();
         const typename range_type::nd_range_type& l_size = work_size.local_range();
         cl_uint s;
         for ( s = 1; s < 3 && g_size[s] != size_t(-1); ++s) {}
         cl_event event;
         enforce_cl_retcode(
             clEnqueueNDRangeKernel( device.my_cl_command_queue, kernel, s,
                 g_offset.data(), g_size.data(), l_size[0] ? l_size.data() : NULL, num_events, num_events ? event_list : NULL, &event ),
             "Failed to enqueue a kernel" );
         return event;
     }

template<typename DeviceFilter>

template<typename T , typename... Rest>

void tbb::flow::interface11::opencl_factory< DeviceFilter >::update_arg_list	(	cl_event	e,
		T &	t,
		Rest &...	args
	)

inlineprivate

Definition at line 958 of file flow_graph_opencl_node.h.

                                                             {
         update_one_arg( e, t );
         update_arg_list( e, args... );
     }

template<typename DeviceFilter>

void tbb::flow::interface11::opencl_factory< DeviceFilter >::update_arg_list ( cl_event )

inlineprivate

Definition at line 963 of file flow_graph_opencl_node.h.

963 {}

template<typename DeviceFilter>

template<typename T >

void tbb::flow::interface11::opencl_factory< DeviceFilter >::update_one_arg	(	cl_event	,
		T &
	)

inlineprivate

Definition at line 950 of file flow_graph_opencl_node.h.

950 {}

template<typename DeviceFilter>

template<typename T , typename F >

void tbb::flow::interface11::opencl_factory< DeviceFilter >::update_one_arg	(	cl_event	e,
		opencl_async_msg< T, F > &	msg
	)

inlineprivate

Definition at line 953 of file flow_graph_opencl_node.h.

                                                                    {
         msg.set_event( e );
     }

Friends And Related Function Documentation

template<typename DeviceFilter>

template<typename Factory >

friend class opencl_buffer_impl

friend

Definition at line 1176 of file flow_graph_opencl_node.h.

template<typename DeviceFilter>

template<typename Factory >

friend class opencl_memory

friend

Definition at line 1178 of file flow_graph_opencl_node.h.

template<typename DeviceFilter>

template<typename Factory >

friend class opencl_program

friend

Definition at line 1174 of file flow_graph_opencl_node.h.

Member Data Documentation

template<typename DeviceFilter>

cl_context tbb::flow::interface11::opencl_factory< DeviceFilter >::my_cl_context

private

Definition at line 1169 of file flow_graph_opencl_node.h.

template<typename DeviceFilter>

opencl_device_list tbb::flow::interface11::opencl_factory< DeviceFilter >::my_devices

private

Definition at line 1168 of file flow_graph_opencl_node.h.

template<typename DeviceFilter>

tbb::spin_mutex tbb::flow::interface11::opencl_factory< DeviceFilter >::my_devices_mutex

private

Definition at line 1171 of file flow_graph_opencl_node.h.

template<typename DeviceFilter>

std::once_flag tbb::flow::interface11::opencl_factory< DeviceFilter >::my_once_flag

private

Definition at line 1167 of file flow_graph_opencl_node.h.

The documentation for this class was generated from the following file:

flow_graph_opencl_node.h

Classes

Public Types

Public Member Functions

Private Member Functions

Static Private Member Functions

Private Attributes

Friends

Detailed Description

template<typename DeviceFilter> class tbb::flow::interface11::opencl_factory< DeviceFilter >

Member Typedef Documentation

Constructor & Destructor Documentation

Member Function Documentation

Friends And Related Function Documentation

Member Data Documentation

template<typename DeviceFilter>
class tbb::flow::interface11::opencl_factory< DeviceFilter >