SYCL_EXTERNAL issue when separating defitions from declarations of a named function object

[PC-1]

Hello, 

I'm new to C++ SYCL. I got the SYCL_EXTERNAL issue:

"error : SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute"

when using a named function object as kernel and trying to separate definitions from declarations in a *.cpp file

The project is built on Visual Studio 2022 under Win11. Intel oneAPI version 2023.0. Below is the code. Hope anyone could help me get the code run.

Thanks so much!

Best regards

PC

"my_class.h" file_________________________________

#ifndef MY_CLASS_H
#define MY_CLASS_H
#pragma once

#include <CL/sycl.hpp>

class my_class {

private:
sycl::accessor<int> inputValues;

public:
//Member functions
//Constructors & Destructors
my_class(const sycl::accessor<int>& in_Val);

my_class() = default;
my_class(const my_class& rhs) = default;
my_class(my_class&& rhs) noexcept = default;
~my_class() = default;

//Assignment
my_class& operator = (const my_class& rhs) = default;
my_class& operator = (my_class&& rhs) noexcept = default;

//Methods
void operator()(sycl::nd_item<1> it, auto& sum) const;
};
#endif

__________________________________________________

"my_class.cpp" file_________________________________

#include "pch.h"

my_class::my_class(const sycl::accessor<int>& in_Val) :inputValues(in_Val) {

};

//Methods
void my_class::operator()(sycl::nd_item<1> it, auto& sum) const {
size_t idx = it.get_global_id(0);
sum += inputValues[idx];
};

________________________________________________________

"main.cpp" file________________________________

int main(){

using namespace std;

size_t N = 1024;
size_t L = 256;

sycl::buffer<int> valuesBuf(N);
{
sycl::host_accessor a(valuesBuf, sycl::write_only);
for (int i = 0; i < N; i++) {
a[i] = i;
}
//destruct host_accessor
}

int sumResult = 0;
sycl::buffer<int> sumBuf(&sumResult, 1);

sycl::queue Q(sycl::gpu_selector_v);

Q.submit([&](sycl::handler& cgh) {

auto inputValues = valuesBuf.get_access(cgh);
auto sumReduction = sycl::reduction(sumBuf, cgh, sycl::plus<>());

my_class m_class(inputValues);
cgh.parallel_for(sycl::nd_range<1> {N, L}, sumReduction, m_class);
});

int sum = numeric_limits<int>::max();
{
sycl::host_accessor sumBuf_hacc(sumBuf, sycl::read_only);
sum = sumBuf_hacc[0];
}

std::cout << "sum is " << sum << endl;
return;

}

[SeshaP_Intel · Accepted solution]

Hi,

Thank you for posting in Intel Communities.

Please find the below modified DPC++ source file which is executing without any errors in Visual Studio 17 2022.

#include <iostream>
#include <CL/sycl.hpp>
#include <numeric>
#include <limits>

using namespace sycl;
using namespace std;

int main()
{
    size_t N = 1024;
    size_t L = 256;

    sycl::buffer<int> valuesBuf(N);
    {
        sycl::host_accessor a(valuesBuf, sycl::write_only);
        for (int i = 0; i < N; i++)
        {
            a[i] = i;
        }
        //destruct host_accessor
    }

    int sumResult = 0;
    sycl::buffer<int> sumBuf(&sumResult, 1);
    sycl::queue Q(sycl::gpu_selector_v);


    Q.submit([&](sycl::handler& cgh) {

        auto inputValues = valuesBuf.get_access(cgh);
    auto sumReduction = sycl::reduction(sumBuf, cgh, sycl::plus<>());

    cgh.parallel_for(sycl::nd_range<1> {N, L}, sumReduction, [=](nd_item<1> it, auto& sum) {
        size_t idx = it.get_global_id(0);
    sum += inputValues[idx];
        });
        });


    int sum = numeric_limits<int>::max();
    {
        sycl::host_accessor sumBuf_hacc(sumBuf, sycl::read_only);
        sum = sumBuf_hacc[0];
    }

    std::cout << "sum is " << sum << "\n";

    return 0;
}

If you still facing the same error, try to build the project in Release(x64) mode.  

I can observe that you are trying to call the function inside the kernel that is defined in a different translational unit. If you want to define a function then you need to make sure that the function is present in the same translation unit. Those functions need to be labeled with SYCL_EXTERNAL.

Please refer to the Data Parallel C++ textbook by James Reinders (Chapter 13) for more information. 

Thanks and Regards,

Pendyala Sesha Srinivas

[SeshaP_Intel]

Hi,

Thanks for accepting our solution. If you need any additional information, please post a new question as this thread will no longer be monitored by Intel.

Thanks and Regards,

Pendyala Sesha Srinivas