SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute

Question

I am trying to calculate the euclidean distance for KNN but in parallel using dpc++. the training dataset contains 5 features and 1600 rows, while I want to calculate the distance between the current test point and each training point on the grid in parallel, but I keep getting an error regarding sycl kernal. code for the function:

code

std::vector<double> distance_calculation_FPGA(queue& q,const std::vector<std::vector<double>>& dataset,const std::vector<double>& curr_test) {
range<1> num_items{ dataset.size()};
std::vector<double>res;


res.resize(dataset.size());
buffer dataset_buf(dataset);
buffer curr_test_buf(curr_test);
buffer res_buf(res.data(), num_items);

q.submit([&](handler& h) {
    accessor a(dataset_buf, h, read_only);
    accessor b(curr_test_buf, h, read_only);

    accessor dif(res_buf, h, write_only, no_init);

   h.parallel_for(num_items, [=](auto i) {
  
        for (int j = 0; j <(const int) a[i].size(); ++j) {
            dif[i] += (a[i][j] - b[j]) * (a[i][j] - b[j]) ;
        }
        });
 });
for (int i = 0; i < res.size(); ++i) {
    std::cout << res[i] << std::endl;
} 
//old distance calculation (serial)
//for (int i = 0; i < dataset.size(); ++i) {
 //   double dis = 0;
   // for (int j = 0; j < dataset[i].size(); ++j) {
     //   dis += (curr_test[j] - dataset[i][j]) * (curr_test[j] - dataset[i][j]);
    //}
    //res.push_back(dis);
//}

return res;

}

the error I am receiving:

SYCL kernel cannot call a variadic function

SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute

Would be extremely grateful for any help!

Thanks

Answer 1

We tried running your code by creating dummy 'dataset' and 'curr_test' variables. We were able to run the program successfully. Please refer this thread

Please refer to the complete code attached below.

#include <CL/sycl.hpp>
#include <iostream> 
using namespace sycl;

std::vector<double> distance_calculation_FPGA(queue& q,const std::vector<std::vector<double>>& dataset,const std::vector<double>& curr_test) 
{
range<1> num_items{ dataset.size()};
std::vector<double>res;

res.resize(dataset.size());
buffer dataset_buf(dataset);
buffer curr_test_buf(curr_test);
buffer res_buf(res.data(), num_items);

q.submit([&](handler& h) {
accessor a(dataset_buf, h, read_only);
accessor b(curr_test_buf, h, read_only);

accessor dif(res_buf, h, write_only, no_init);

h.parallel_for(num_items, [=](auto i) {

    for (int j = 0; j <(const int) a[i].size(); ++j) {
//           dif[i] += (a[i][j] - b[j]) * (a[i][j] - b[j]) ;
         dif[i]+=a[i][j];          
        }
        });
    });
    q.wait(); //We have added this line of code for synchronization.
    for (int i : res) { 
        std::cout <<i<< std::endl;
    } 
    return res;
    }


int main(){
std::vector<std::vector<double>> dataset;
for(int i=0;i<5;i++)
{
    std::vector<double> d;

    for(int j=0;j<1600;j++)
    {
        d.push_back((double)j);
    }
    dataset.push_back(d);
}

std::vector<double> curr_test;
for(int i=0;i<1600;i++)
{
    curr_test.push_back((double)i);
}
queue q;
std::cout << "Running on "<< 
q.get_device().get_info<sycl::info::device::name>()<< std::endl; 
//print the device name as a test to check the parallelisation

distance_calculation_FPGA(q,dataset,curr_test);

return 0;
}