Intel® oneAPI Data Parallel C++
Support for Intel® oneAPI DPC++ Compiler, Intel® oneAPI DPC++ Library, Intel® DPC++ Compatibility Tool, and GDB*
547 Discussions

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

PC-1
Beginner
430 Views

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;

}

 

0 Kudos
1 Solution
SeshaP_Intel
Moderator
386 Views

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

 

View solution in original post

0 Kudos
2 Replies
SeshaP_Intel
Moderator
387 Views

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

 

0 Kudos
SeshaP_Intel
Moderator
361 Views

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


0 Kudos
Reply