CLTuner Class Reference

Basic implementation of the OpenCL tuner interface. More...

#include <CLTuner.h>

Collaboration diagram for CLTuner:

Public Member Functions
	CLTuner (bool tune_new_kernels=true)
	Constructor. More...
	~CLTuner ()=default
	Destructor. More...
void	set_tune_new_kernels (bool tune_new_kernels)
	Setter for tune_new_kernels option. More...
bool	tune_new_kernels () const
	Tune kernels that are not in the LWS table. More...
void	set_tuner_mode (CLTunerMode mode)
	Set OpenCL tuner mode. More...
CLTunerMode	get_tuner_mode () const
	Get the current OpenCL tuner mode. More...
void	add_lws_to_table (const std::string &kernel_id, cl::NDRange optimal_lws)
	Manually add a LWS for a kernel. More...
void	import_lws_table (const std::unordered_map< std::string, cl::NDRange > &lws_table)
	Import LWS table. More...
const std::unordered_map< std::string, cl::NDRange > &	lws_table () const
	Give read access to the LWS table. More...
void	set_cl_kernel_event (cl_event kernel_event)
	Set the OpenCL kernel event. More...
void	load_from_file (const std::string &filename)
	Load the LWS table from file. More...
void	save_to_file (const std::string &filename) const
	Save the content of the LWS table to file. More...
void	tune_kernel_static (ICLKernel &kernel) override
	Tune OpenCL kernel statically. More...
void	tune_kernel_dynamic (ICLKernel &kernel) override
	Tune OpenCL kernel dynamically. More...
bool	kernel_event_is_set () const
	Is the kernel_event set ? More...
Public Member Functions inherited from ICLTuner
virtual	~ICLTuner ()=default
	Virtual destructor. More...

Data Fields
std::function< decltype(clEnqueueNDRangeKernel)>	real_clEnqueueNDRangeKernel
	clEnqueueNDRangeKernel symbol More...

Detailed Description

Basic implementation of the OpenCL tuner interface.

Definition at line 38 of file CLTuner.h.

Constructor & Destructor Documentation

CLTuner()

CLTuner

(

bool

tune_new_kernels = true

)

Constructor.

Parameters

[in]

tune_new_kernels

Find the optimal local workgroup size for kernels which are not present in the table ?

Definition at line 40 of file CLTuner.cpp.

    : real_clEnqueueNDRangeKernel(nullptr), _lws_table(), _kernel_event(), _tune_new_kernels(tune_new_kernels), _tuner_mode(CLTunerMode::NORMAL)
{
}

References arm_compute::NORMAL.

~CLTuner()

~CLTuner ( )

default

Destructor.

Member Function Documentation

add_lws_to_table()

void add_lws_to_table	(	const std::string &	kernel_id,
		cl::NDRange	optimal_lws
	)

Manually add a LWS for a kernel.

Parameters

[in]	kernel_id	Unique identifiant of the kernel
[in]	optimal_lws	Optimal local workgroup size to use for the given kernel

Definition at line 109 of file CLTuner.cpp.

{
    _lws_table.emplace(kernel_id, optimal_lws);
}

Referenced by CLTuner::tune_kernel_dynamic().

get_tuner_mode()

CLTunerMode get_tuner_mode

(

)

const

Get the current OpenCL tuner mode.

Returns

tuner_mode Indicates how exhaustive the search for the optimal LWS should be while tuning

Definition at line 67 of file CLTuner.cpp.

{
    return _tuner_mode;
}

import_lws_table()

void import_lws_table

(

const std::unordered_map< std::string, cl::NDRange > &

lws_table

)

Import LWS table.

Parameters

[in]

lws_table

The unordered_map container to import

Definition at line 222 of file CLTuner.cpp.

{
    _lws_table.clear();
    _lws_table = lws_table;
}

References CLTuner::lws_table().

kernel_event_is_set()

bool kernel_event_is_set

(

)

const

Is the kernel_event set ?

Returns

true if the kernel_event is set.

Definition at line 45 of file CLTuner.cpp.

{
    return _kernel_event() != nullptr;
}

load_from_file()

void load_from_file

(

const std::string &

filename

)

Load the LWS table from file.

Parameters

[in]

filename

Load the LWS table from this file.(Must exist)

Definition at line 233 of file CLTuner.cpp.

{
    std::ifstream fs;
    fs.exceptions(std::ifstream::badbit);
    fs.open(filename, std::ios::in);
    if(!fs.is_open())
    {
        ARM_COMPUTE_ERROR("Failed to open '%s' (%s [%d])", filename.c_str(), strerror(errno), errno);
    }
    std::string line;
    while(!std::getline(fs, line).fail())
    {
        std::istringstream ss(line);
        std::string        token;
        if(std::getline(ss, token, ';').fail())
        {
            ARM_COMPUTE_ERROR("Malformed row '%s' in %s (Should be of the form 'kernel_id;lws[0];lws[1];lws[2]')", ss.str().c_str(), filename.c_str());
        }
        std::string kernel_id = token;
        cl::NDRange lws(1, 1, 1);
        for(int i = 0; i < 3; i++)
        {
            if(std::getline(ss, token, ';').fail())
            {
                ARM_COMPUTE_ERROR("Malformed row '%s' in %s (Should be of the form 'kernel_id;lws[0];lws[1];lws[2]')", ss.str().c_str(), filename.c_str());
            }
            lws.get()[i] = support::cpp11::stoi(token);
        }

        // If all dimensions are 0: reset to NullRange (i.e nullptr)
        if(lws[0] == 0 && lws[1] == 0 && lws[2] == 0)
        {
            lws = cl::NullRange;
        }
        add_lws_to_table(kernel_id, lws);
    }
    fs.close();
}

References ARM_COMPUTE_ERROR, and arm_compute::support::cpp11::stoi().

Referenced by main(), and CLDeviceBackend::setup_backend_context().

lws_table()

const std::unordered_map< std::string, cl::NDRange > & lws_table

(

)

const

Give read access to the LWS table.

Returns

The lws table as unordered_map container

Definition at line 228 of file CLTuner.cpp.

{
    return _lws_table;
}

Referenced by CLTuner::import_lws_table(), and CLDeviceBackend::~CLDeviceBackend().

save_to_file()

void save_to_file

(

const std::string &

filename

)

const

Save the content of the LWS table to file.

Parameters

[in]

filename

Save the LWS table to this file. (Content will be overwritten)

Definition at line 272 of file CLTuner.cpp.

{
    std::ofstream fs;
    fs.exceptions(std::ifstream::failbit | std::ifstream::badbit);
    fs.open(filename, std::ios::out);
    for(auto const &kernel_data : _lws_table)
    {
        fs << kernel_data.first << ";" << kernel_data.second[0] << ";" << kernel_data.second[1] << ";" << kernel_data.second[2] << std::endl;
    }
    fs.close();
}

Referenced by main(), and CLDeviceBackend::~CLDeviceBackend().

set_cl_kernel_event()

void set_cl_kernel_event

(

cl_event

kernel_event

)

Set the OpenCL kernel event.

Note

The interceptor can use this function to store the event associated to the OpenCL kernel

Parameters

[in]

kernel_event

The OpenCL kernel event

Definition at line 49 of file CLTuner.cpp.

{
    _kernel_event = kernel_event;
}

set_tune_new_kernels()

void set_tune_new_kernels

(

bool

tune_new_kernels

)

Setter for tune_new_kernels option.

Parameters

[in]

tune_new_kernels

Find the optimal local workgroup size for kernels which are not present in the table ?

Definition at line 54 of file CLTuner.cpp.

{
    _tune_new_kernels = tune_new_kernels;
}

References CLTuner::tune_new_kernels().

Referenced by main(), and CLDeviceBackend::set_kernel_tuning().

set_tuner_mode()

void set_tuner_mode

(

CLTunerMode

mode

)

Set OpenCL tuner mode.

Parameters

[in]

mode

Indicates how exhaustive the search for the optimal LWS should be while tuning. Default is Exhaustive mode

Definition at line 63 of file CLTuner.cpp.

{
    _tuner_mode = mode;
}

Referenced by main(), and CLDeviceBackend::set_kernel_tuning_mode().

tune_kernel_dynamic()

void tune_kernel_dynamic ( ICLKernel &  kernel )

overridevirtual

Tune OpenCL kernel dynamically.

Note

Tuning requires memory to be available on all kernel tensors and objects in order to be performed

Parameters

[in]

kernel

Kernel to tune

Implements ICLTuner.

Definition at line 77 of file CLTuner.cpp.

{
    // Get the configuration ID from the kernel and append GPU target name and number of available compute units
    const std::string config_id = kernel.config_id() + "_" + string_from_target(kernel.get_target()) + "_MP" + support::cpp11::to_string(CLKernelLibrary::get().get_num_compute_units());

    // Check if we need to find the Optimal LWS. If the kernel's config_id is equal to default_config_id, the kernel does not require to be tuned
    if(kernel.config_id() != arm_compute::default_config_id)
    {
        auto p = _lws_table.find(config_id);

        if(p == _lws_table.end())
        {
            if(_tune_new_kernels)
            {
                // Find the optimal LWS for the kernel
                cl::NDRange opt_lws = find_optimal_lws(kernel);

                // Insert the optimal LWS in the table
                add_lws_to_table(config_id, opt_lws);

                // Set Local-Workgroup-Size
                kernel.set_lws_hint(opt_lws);
            }
        }
        else
        {
            // Set Local-Workgroup-Size
            kernel.set_lws_hint(p->second);
        }
    }
}

References CLTuner::add_lws_to_table(), ICLKernel::config_id(), CLKernelLibrary::get(), CLKernelLibrary::get_num_compute_units(), ICLKernel::get_target(), ICLKernel::set_lws_hint(), arm_compute::string_from_target(), and arm_compute::support::cpp11::to_string().

tune_kernel_static()

void tune_kernel_static ( ICLKernel &  kernel )

overridevirtual

Tune OpenCL kernel statically.

Note

Tuning is performed using only kernel and tensor metadata, thus can be performed when memory is not available

Parameters

[in]

kernel

Kernel to tune

Implements ICLTuner.

Definition at line 72 of file CLTuner.cpp.

{
    ARM_COMPUTE_UNUSED(kernel);
}

References ARM_COMPUTE_UNUSED.

tune_new_kernels()

bool tune_new_kernels

(

)

const

Tune kernels that are not in the LWS table.

Returns

True if tuning of new kernels is enabled.

Definition at line 58 of file CLTuner.cpp.

{
    return _tune_new_kernels;
}

Referenced by CLTuner::set_tune_new_kernels(), and CLDeviceBackend::~CLDeviceBackend().

Field Documentation

real_clEnqueueNDRangeKernel

std::function<decltype(clEnqueueNDRangeKernel)> real_clEnqueueNDRangeKernel

clEnqueueNDRangeKernel symbol

Definition at line 102 of file CLTuner.h.

---

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

• arm_compute/runtime/CL/CLTuner.h
• src/runtime/CL/CLTuner.cpp