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, CLTuningInfo tuning_info=CLTuningInfo())
	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 tuning parameters table. More...
void	set_tuning_parameters (CLTuningInfo tuning_info)
	Setter for tune parameters option. 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	add_tuning_params (const std::string &kernel_id, CLTuningParams optimal_tuning_params)
	Manually add tuning parameters for a kernel. More...
void	import_lws_table (const std::unordered_map< std::string, cl::NDRange > &lws_table)
	Import LWS table. More...
void	import_tuning_params (const std::unordered_map< std::string, CLTuningParams > &tuning_params_table)
	Import tuning parameters table. More...
const std::unordered_map< std::string, cl::NDRange > &	lws_table ()
	Give read access to the LWS table. More...
const std::unordered_map< std::string, CLTuningParams > &	tuning_params_table () const
	Give read access to the tuning params 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 tuning parameters table from file. More...
bool	save_to_file (const std::string &filename) const
	Save the content of the tuning parameters 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...
void	tune_kernel_dynamic (ICLKernel &kernel, ITensorPack &tensors) 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 40 of file CLTuner.h.

Constructor & Destructor Documentation

CLTuner()

CLTuner	(	bool	tune_new_kernels = true,
		CLTuningInfo	tuning_info = CLTuningInfo()
	)

Constructor.

Parameters

[in]	tune_new_kernels	Find the optimal local workgroup size for kernels which are not present in the table ?
[in]	tuning_info	(Optional) opencl parameters to tune

Definition at line 38 of file CLTuner.cpp.

    : real_clEnqueueNDRangeKernel(nullptr), _tuning_params_table(), _lws_table(), _kernel_event(), _tune_new_kernels(tune_new_kernels), _tuning_info(tuning_info)
{
}

~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

Deprecated:

This function is deprecated and is intended to be removed in 21.05 release

Definition at line 122 of file CLTuner.cpp.

References CLTuner::add_tuning_params().

{
    add_tuning_params(kernel_id, CLTuningParams(optimal_lws));
}

add_tuning_params()

void add_tuning_params	(	const std::string &	kernel_id,
		CLTuningParams	optimal_tuning_params
	)

Manually add tuning parameters for a kernel.

Parameters

[in]	kernel_id	Unique identifiant of the kernel
[in]	optimal_tuning_params	Optimal tuning parameters to use for the given kernel

Definition at line 127 of file CLTuner.cpp.

References CLSymbols::clEnqueueNDRangeKernel_ptr, clRetainEvent(), ITensorPack::empty(), arm_compute::mlgo::parser::end(), CLScheduler::get(), CLSymbols::get(), ICLKernel::gws_from_window(), CLTuner::kernel_event_is_set(), CLScheduler::queue(), CLTuner::real_clEnqueueNDRangeKernel, CLTuner::set_cl_kernel_event(), and enable_tracing::start.

Referenced by CLTuner::add_lws_to_table(), CLTuner::import_lws_table(), CLTuner::load_from_file(), and CLTuner::tune_kernel_dynamic().

{
    _tuning_params_table.emplace(kernel_id, optimal_tuning_params);
}

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 tuning parameters should be while tuning

Deprecated:

This function is deprecated and is intended to be removed in 21.05 release

Definition at line 66 of file CLTuner.cpp.

References CLTuningInfo::tuner_mode.

{
    return _tuning_info.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

Deprecated:

This function is deprecated and is intended to be removed in 21.05 release

Definition at line 250 of file CLTuner.cpp.

References CLTuner::add_tuning_params().

{
    _tuning_params_table.clear();
    for(auto && params : lws_table)
    {
        add_tuning_params(params.first, CLTuningParams(params.second));
    }
}

import_tuning_params()

void import_tuning_params

(

const std::unordered_map< std::string, CLTuningParams > &

tuning_params_table

)

Import tuning parameters table.

Parameters

[in]

tuning_params_table

The unordered_map container to import

Definition at line 274 of file CLTuner.cpp.

References CLTuner::tuning_params_table().

{
    _tuning_params_table.clear();
    _tuning_params_table = tuning_params_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 43 of file CLTuner.cpp.

Referenced by CLTuner::add_tuning_params().

{
    return _kernel_event() != nullptr;
}

load_from_file()

void load_from_file

(

const std::string &

filename

)

Load the tuning parameters table from file.

It also sets up the tuning read from the file

Parameters

[in]

filename

Load the tuning parameters table from this file.(Must exist)

Definition at line 280 of file CLTuner.cpp.

References CLTuner::add_tuning_params(), ARM_COMPUTE_ERROR_VAR, CLTuningParams::from_string(), and CLTuningInfo::tune_wbsm.

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

{
    std::ifstream fs;
    fs.exceptions(std::ifstream::badbit);
    fs.open(filename, std::ios::in);
    if(!fs.is_open())
    {
        ARM_COMPUTE_ERROR_VAR("Failed to open '%s' (%s [%d])", filename.c_str(), strerror(errno), errno);
    }
    std::string line;
    bool        header_line = true;
    while(!std::getline(fs, line).fail())
    {
        if(header_line)
        {
            header_line            = false;
            size_t pos_lws         = line.find("lws");
            size_t pos_wbsm        = line.find("wbsm");
            _tuning_info.tune_wbsm = false;
            if(pos_lws != std::string::npos || pos_wbsm != std::string::npos)
            {
                // The file has in the first line the parameters it has been tuned on
                if(pos_wbsm != std::string::npos)
                {
                    _tuning_info.tune_wbsm = true;
                }
                // Once the line with the tuning parameter is read we can
                // read the next one to start collecting the values
                if(std::getline(fs, line).fail())
                {
                    break;
                }
            }
        }

        CLTuningParams tuning_params;
        size_t         pos = line.find(";");
        if(pos == std::string::npos)
        {
            ARM_COMPUTE_ERROR_VAR("Malformed row '%s' in %s", line.c_str(), filename.c_str());
        }
        std::string kernel_id = line.substr(0, pos);
        line.erase(0, pos + 1);
        if(!tuning_params.from_string(_tuning_info, line))
        {
            ARM_COMPUTE_ERROR_VAR("Malformed row '%s' in %s", line.c_str(), filename.c_str());
        }
        add_tuning_params(kernel_id, tuning_params);
    }
    fs.close();
}

lws_table()

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

(

)

Give read access to the LWS table.

Returns

The lws table as unordered_map container

Deprecated:

This function is deprecated and is intended to be removed in 21.05 release

Definition at line 259 of file CLTuner.cpp.

{
    _lws_table.clear();
    for(auto && params : _tuning_params_table)
    {
        _lws_table.emplace(params.first, params.second.get_lws());
    }
    return _lws_table;
}

save_to_file()

bool save_to_file

(

const std::string &

filename

)

const

Save the content of the tuning parameters table to file.

Parameters

[in]

filename

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

Returns

true if the file was created

Definition at line 332 of file CLTuner.cpp.

References CLTuningParams::to_string(), and CLTuningInfo::tune_wbsm.

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

{
    if(!_tune_new_kernels || _tuning_params_table.empty() || filename.empty())
    {
        return false;
    }
    std::ofstream fs;
    fs.exceptions(std::ifstream::failbit | std::ifstream::badbit);
    fs.open(filename, std::ios::out);
    std::string header_string = "";
    header_string += "lws";
    if(_tuning_info.tune_wbsm)
    {
        if(!header_string.empty())
        {
            header_string += " ";
        }
        header_string += "wbsm";
    }
    fs << header_string << std::endl;
    for(auto const &kernel_data : _tuning_params_table)
    {
        CLTuningParams tun_pams(kernel_data.second);
        fs << kernel_data.first << tun_pams.to_string(_tuning_info) << std::endl;
    }
    fs.close();
    return true;
}

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 47 of file CLTuner.cpp.

Referenced by CLTuner::add_tuning_params().

{
    _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 52 of file CLTuner.cpp.

References CLTuner::tune_new_kernels().

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

{
    _tune_new_kernels = tune_new_kernels;
}

set_tuner_mode()

void set_tuner_mode

(

CLTunerMode

mode

)

Set OpenCL tuner mode.

Parameters

[in]

mode

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

Definition at line 61 of file CLTuner.cpp.

References clang_tidy_rules::mode, and CLTuningInfo::tuner_mode.

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

{
    _tuning_info.tuner_mode = mode;
}

set_tuning_parameters()

void set_tuning_parameters

(

CLTuningInfo

tuning_info

)

Setter for tune parameters option.

Parameters

[in]

tuning_info

opencl parameters to tune

tune_kernel_dynamic() [1/2]

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 76 of file CLTuner.cpp.

{
    ITensorPack pack;
    tune_kernel_dynamic(kernel, pack);
}

tune_kernel_dynamic() [2/2]

void tune_kernel_dynamic ( ICLKernel &  kernel, ITensorPack &  tensors  )

overridevirtual

Tune OpenCL kernel dynamically.

Parameters

[in]	kernel	Kernel to tune
[in,out]	tensors	Tensors for the kernel to use

Implements ICLTuner.

Definition at line 82 of file CLTuner.cpp.

References CLTuner::add_tuning_params(), ICLKernel::config_id(), CLKernelLibrary::get(), CLTuningParams::get_lws(), CLKernelLibrary::get_num_compute_units(), ICLKernel::get_target(), CLTuningParams::get_wbsm(), ICLKernel::set_lws_hint(), ICLKernel::set_wbsm_hint(), arm_compute::string_from_target(), arm_compute::support::cpp11::to_string(), and CLTuningInfo::tune_wbsm.

{
    // 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 = _tuning_params_table.find(config_id);

        if(p == _tuning_params_table.end())
        {
            if(_tune_new_kernels)
            {
                // Find the optimal LWS for the kernel
                CLTuningParams opt_tuning_params = find_optimal_tuning_params(kernel, tensors);

                // Insert the optimal LWS in the table
                add_tuning_params(config_id, opt_tuning_params);

                // Set Local-Workgroup-Size
                kernel.set_lws_hint(opt_tuning_params.get_lws());
                if(_tuning_info.tune_wbsm)
                {
                    kernel.set_wbsm_hint(opt_tuning_params.get_wbsm());
                }
            }
        }
        else
        {
            // Set Local-Workgroup-Size
            kernel.set_lws_hint(p->second.get_lws());
            if(_tuning_info.tune_wbsm)
            {
                kernel.set_wbsm_hint(p->second.get_wbsm());
            }
        }
    }
}

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 71 of file CLTuner.cpp.

References ARM_COMPUTE_UNUSED.

{
    ARM_COMPUTE_UNUSED(kernel);
}

tune_new_kernels()

bool tune_new_kernels

(

)

const

Tune kernels that are not in the tuning parameters table.

Returns

True if tuning of new kernels is enabled.

Definition at line 56 of file CLTuner.cpp.

Referenced by CLTuner::set_tune_new_kernels().

{
    return _tune_new_kernels;
}

tuning_params_table()

const std::unordered_map< std::string, CLTuningParams > & tuning_params_table

(

)

const

Give read access to the tuning params table.

Returns

The tuning params table as unordered_map container

Definition at line 269 of file CLTuner.cpp.

Referenced by CLTuner::import_tuning_params().

{
    return _tuning_params_table;
}

Field Documentation

real_clEnqueueNDRangeKernel

std::function<decltype(clEnqueueNDRangeKernel)> real_clEnqueueNDRangeKernel

clEnqueueNDRangeKernel symbol

Definition at line 142 of file CLTuner.h.

Referenced by CLTuner::add_tuning_params().

---

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

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