Task Identification with KeyArgs

The template parameter pack KeyArgs allows the policy to track performance for submitted jobs. By default (empty KeyArgs), all invocations of the same function share performance history. When KeyArgs are specified, the policy uses both the function pointer and the specified arguments to create a unique key for tracking performance.

For example, auto_tune_policy<sycl::queue, oneapi::dpl::identity, default_backend, std::size_t> will track performance separately for each distinct std::size_t argument value, useful when performance varies with problem size. Here, one KeyArg type corresponds to one extra argument passed to the function after the sycl::queue.

Example

In the following example, an auto_tune_policy is used to dynamically select between two queues, a CPU queue and a GPU queue.

#include <oneapi/dpl/dynamic_selection>
#include <sycl/sycl.hpp>
#include <iostream>

namespace ex = oneapi::dpl::experimental;

int main() {
  std::vector<sycl::queue> r { sycl::queue{sycl::cpu_selector_v},
                               sycl::queue{sycl::gpu_selector_v} };

  const std::size_t N = 10000;
  std::vector<float> av(N, 0.0);
  std::vector<float> bv(N, 0.0);
  std::vector<float> cv(N, 0.0);
  for (int i = 0; i < N; ++i) {
    av[i] = bv[i] = i;
  }

  ex::auto_tune_policy p{r}; // (1)

  {
    sycl::buffer<float> a_b(av);
    sycl::buffer<float> b_b(bv);
    sycl::buffer<float> c_b(cv);


    for (int i = 0; i < 6; ++i) {
      ex::submit_and_wait(p, [&](sycl::queue q) { // (2)
        // (3)
        std::cout << (q.get_device().is_cpu() ? "using cpu\n" : "using gpu\n");
        return q.submit([&](sycl::handler &h) { // (4)
          sycl::accessor a_a(a_b, h, sycl::read_only);
          sycl::accessor b_a(b_b, h, sycl::read_only);
          sycl::accessor c_a(c_b, h, sycl::read_write);
          h.parallel_for(N, [=](auto i) { c_a[i] = a_a[i] + b_a[i]; });
        });
      });
    };
  }

  for (int i = 0; i < N; ++i) {
    if (cv[i] != 2*i) {
       std::cout << "ERROR!\n";
    }
  }
  std::cout << "Done.\n";
}

The key points in this example are:

1. An auto_tune_policy is constructed to select between the CPU and GPU.

2. submit_and_wait is invoked with the policy as the first argument. The selected queue will be passed to the user-provided function.

3. For clarity when run, the type of device is displayed.

4. The queue is used in function to perform and asynchronous offload. The SYCL event returned from the call to submit is returned. Returning an event is required for functions passed to submit and submit_and_wait.

Selection Algorithm

The selection algorithm for auto_tune_policy uses runtime profiling to choose the best resource for the given function. A simplified, expository implementation of the selection algorithm follows:

//not a public function, for exposition purposes only
template<typename Function, typename ...Args>
selection_type auto_tune_policy::select(Function&& f, Args&&...args) {
  if (initialized_) {
    auto k = make_task_key(f, args...);
    auto tuner = get_tuner(k);
    auto offset = tuner->get_resource_to_profile();
    if (offset == use_best) {
      return selection_type{*this, tuner->best_resource_, tuner};
    } else {
      auto r = resources_[offset];
      return selection{*this, r, tuner};
    }
  } else {
    throw std::logic_error("selected called before initialization");
  }
}

where make_task_key combines the inputs, including the function and its arguments, into a key that uniquely identifies the user function that is being profiled. tuner is the encapsulated logic for performing runtime profiling and choosing the best option for a given key. When the call to get_resource_to_profile() return use_best, the tuner is not in the profiling phase, and so the previously determined best resource is used. Otherwise, the resource at index offset in the resources_ vector is used and its resulting performance is profiled. When an auto_tune_policy is initialized with a non-zero resample interval, the policy will periodically return to the profiling phase base on the provided interval value.

Constructors

auto_tune_policy provides three constructors.

Deferred Initialization

An auto_tune_policy that was constructed with deferred initialization must be initialized by calling one of its initialize member functions before it can be used to select or submit.

Queries

An auto_tune_policy has get_resources and get_submission_group member functions.

Reporting Requirements

A auto_tune_policy requires the task_time reporting requirement. See the Execution Information section for more information about reporting requirements.