How to define and execute an array of functions on Sycl+openCL+DPCPP

Question

In my program, I defined an array of functions

    #include <CL/sycl.hpp>
#include <iostream>
#include <tbb/tbb.h>
#include <tbb/parallel_for.h>
#include <vector>
#include <string>
#include <queue>
#include<tbb/blocked_range.h>
#include <tbb/global_control.h>
#include <chrono>


using namespace tbb;

template<class Tin, class Tout, class Function>
class Map {
private:
    Function fun;
public:
    Map() {}
    Map(Function f):fun(f) {}


    std::vector<Tout> operator()(bool use_tbb, std::vector<Tin>& v) {
        std::vector<Tout> r(v.size());
        if(use_tbb){
            // Start measuring time
            auto begin = std::chrono::high_resolution_clock::now();
            tbb::parallel_for(tbb::blocked_range<Tin>(0, v.size()),
                        [&](tbb::blocked_range<Tin> t) {
                    for (int index = t.begin(); index < t.end(); ++index){
                        r[index] = fun(v[index]);
                    }
            });
            // Stop measuring time and calculate the elapsed time
            auto end = std::chrono::high_resolution_clock::now();
            auto elapsed = std::chrono::duration_cast<std::chrono::nanoseconds>(end - begin);
            printf("Time measured: %.3f seconds.\n", elapsed.count() * 1e-9);
            return r;
         } else {
                sycl::queue gpuQueue{sycl::gpu_selector()};
                sycl::range<1> n_item{v.size()};
                sycl::buffer<Tin, 1> in_buffer(&v[0], n_item);
                sycl::buffer<Tout, 1> out_buffer(&r[0], n_item);
                gpuQueue.submit([&](sycl::handler& h){
                    //local copy of fun
                    auto f = fun;
                    sycl::accessor in_accessor(in_buffer, h, sycl::read_only);
                    sycl::accessor out_accessor(out_buffer, h, sycl::write_only);
                    h.parallel_for(n_item, [=](sycl::id<1> index) {
                        out_accessor[index] = f(in_accessor[index]);
                    });
                }).wait();
         }
                return r;
    }
};

template<class Tin, class Tout, class Function>
Map<Tin, Tout, Function> make_map(Function f) { return Map<Tin, Tout, Function>(f);}


typedef int(*func)(int x);
//define different functions
auto function = [](int x){ return x; };
auto functionTimesTwo = [](int x){ return (x*2); };
auto functionDivideByTwo = [](int x){ return (x/2); };
auto lambdaFunction = [](int x){return (++x);};


int main(int argc, char *argv[]) {

    std::vector<int> v = {1,2,3,4,5,6,7,8,9};
    //auto f = [](int x){return (++x);};
    //Array of functions
    func functions[] =
        {
            function,
            functionTimesTwo,
            functionDivideByTwo,
            lambdaFunction
        };

    for(int i = 0; i< sizeof(functions); i++){
        auto m1 = make_map<int, int>(functions[i]);

    //auto m1 = make_map<int, int>(f);
    std::vector<int> r = m1(true, v);
    //print the result
    for(auto &e:r) {
        std::cout << e << " ";
        }
    }


  return 0;
}

instead of each time defining a function, I am interested in defining an array of functions and then execute it in my program. But in the part of SYCL for executing on GPU, I have an error and I do not know how to fix it.

The ERROR: SYCL kernel cannot call through a function pointer

Answer 1

In particular, SYCL device code, as defined by this specification, does not support virtual function calls, function pointers in general, exceptions, runtime type information or the full set of C++ libraries that may depend on these features or on features of a particular host compiler. Nevertheless, these basic restrictions can be relieved by some specific Khronos or vendor extensions.

As per the sycl 2020 specification, No function pointers are allowed to be called in a SYCL kernel or any functions called by the kernel. Please refer https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#introduction