Calling function from kernel scope. SYCL_EXTERNAL

[Xeenych]

I have my_func.cpp/.hpp with some fuction

/* my_func.hpp */
void my_func();

/* my_func.cpp */
void my_func() {
// do some stuff   
// call some other funcs
}

In my main.cpp I wrote a simple kernel, which just calls my_func()

When compiling it with VS2022 compiler says, that my_func() should have SYCL_EXTERNAL attribute?

How do i get over this error?

[NoorjahanSk_Intel]

Hi,

Thanks for reaching out to us.

>>How do i get over this error?

SYCL_EXTERNAL is an optional macro that enables external linkage of SYCL functions and methods to be included in an SYCL kernel.

If we want to call a function from a kernel, that function(in our case void my_func()) should be preceded with SYCL_EXTERNAL in the header file, then you can get rid of this compiler error.

We have tried the following code with dpc++ version 2022.0.0 and it worked for us.

Please find the below code snippets.

Header.h

#pragma once
#include<CL/sycl.hpp>
extern SYCL_EXTERNAL void my_func();

my_fun.cpp

#include <CL/sycl.hpp>
#include "Header.h"
void my_func() {
}

main.cpp

#include <CL/sycl.hpp>
#include <iostream>
#include "Header.h"
#define numElements 10
using namespace std;

int main(void) 
{
cl::sycl::queue queue(cl::sycl::host_selector{});
std::cout << "Running on " << queue.get_device().get_info<cl::sycl::info::device::name>() << "\n";

{
queue.submit(
[&](cl::sycl::handler& cgh) {
cgh.parallel_for<class vectorAdd_e83213>(
cl::sycl::range<1>{numElements}, [=](cl::sycl::item<1> item_ct1) {
my_func(); 
});
});
}
queue.wait();
return 0;
}

If this resolves your issue, make sure to accept this as a solution. This would help others with a similar issue.

If your issue still persists, please do let us know which oneAPI version is being used.

Thanks & Regards,

Noorjahan

[Xeenych]

Ok, this works, but:

my_func() calls tons of other functions.

I have to add SYCL_EXTERNAL to every declaration in .hpp and every definition in.cpp of all functions in a call-tree of my_func()

This introduces a dependancy on DPC++ headers in my_func() and other sources

I could have written something like this for every function in a call-tree of my_func().

#if defined __INTEL_LLVM_COMPILER && SYCL_LANGUAGE_VERISON
#include <CL/sycl.hpp>
#else
#define SYCL_EXTERNAL
#endif

SYCL_EXTERNAL void my_func();

Is there a way to avoid this?

Is there a linker script or some command-line arguments for compiler to list functions to be included in SYCL kernel without specifying SYCL_EXTERNAL?

I thought, that intel compiler is advanced enough to determine functions to be included in kernel and link them accordingly.

[NoorjahanSk_Intel · Accepted solution]

Hi,

>>Is there a linker script or some command-line arguments for compiler to list functions to be included in SYCL kernel without specifying SYCL_EXTERNAL?

Sorry for causing any inconvenience, there is no such script or command-line argument to list the functions.

As of now, we need to use macro definitions to add SYCL_EXTERNAL for every function as you have mentioned in the earlier post.

Thanks & Regards,

Noorjahan.

[NoorjahanSk_Intel]

Hi,

Could you please provide an update on your issue?

If the information provided above helped, Could you also please confirm whether can we close this issue from our end.

Thanks & Regards,

Noorjahan.

[Xeenych]

Yes, this helped. You can close the issue.

P.S.

Looks like my account was banned for several days, but now it works.

[NoorjahanSk_Intel]

Hi,

Thanks for the confirmation!

As this issue has been resolved, we will no longer respond to this thread. If you need any additional information, please submit a new question as this thread will no longer be monitored.

Thanks & Regards,

Noorjahan.