SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute

[Jeremy_Wang]

Hi, I'm trying to use oneAPI to calculate an index. The dataset has 6 cols and 2700 rows. While I want to calculate the index, I keep getting an error with sycl kernel. code for the function:

//Initialize data from local csv file
    vector<vector<float>> data = read("..\\data\\dataset.csv");
    vector<float> weight(data[1].size());
    vector<float> MEI(data.size());
    vector<float> init = data[0];
    
    buffer buf_data(data);
    buffer buf_init(init);
    buffer<float, 2> buf_temp_MEI{ range{data.size(), data[1].size()} };
    buffer<float, 1> buf_weight(weight.data(), weight.size());
    buffer<float, 1> buf_MEI(MEI.data(), MEI.size());
    
    try {
    
        queue q(d_selector, exception_handler);
        // Print out the device information used for the kernel code.
        std::cout << "Running on device: "
            << q.get_device().get_info<info::device::name>() << "\n";

        //Calculate Weight for each assets
        q.submit([&](handler& h) {
            accessor ass_init(buf_init, h, read_only);
            // The sum_accessor is used to store (with write permission) the sum data.
            accessor ass_weight(buf_weight, h, read_write, noinit);
            h.parallel_for(weight.size(), [=](auto i) { ass_weight[i] = (100.0 / ass_init[i]); });
            
        });
        
        
        host_accessor ass(buf_temp_MEI, read_only);
        for (int i = 0; i < weight.size(); i++) {
            cout << "Weight for " << i << ": " << ass[0][i] << std::endl;
        }
        
        
        q.submit([&](handler& h) {
            accessor ass_weight(buf_weight, h, read_only);
            accessor ass_data(buf_data, h, read_only);
            accessor ass_temp(buf_temp_MEI, h, write_only, noinit);
            h.parallel_for(sycl::range<2>(data.size(), data[0].size()),
                [=](sycl::id<2> index) {
                    int row = index[0];
                    int col = index[1];
                    ass_temp[index] = ass_data[row][col] * ass_weight[col];
                });
            
        });
    }
    catch (sycl::exception const& e) {
        std::cout << "An exception is caught.\n";
        std::terminate();
    }

    return 0;

I test all the codes inside the program. And I'm sure there is something wrong with the second q.submit(), which is the one having 2-d parallel for function. 

The error I'm receving: 

SYCL kernel cannot call a variadic function
SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute
SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute

Would be extremely grateful for any help!

Thanks

[VidyalathaB_Intel]

Hi Jeremy,

Thanks for reaching out to us.

We tried to reproduce the error that you are getting but could not see the error. So could you please provide us with the following details so that we could try to address the issue that you are getting?

The version of the DPC++ compiler being used (command to check: dpcpp --version)

OS environment details

Please attach the Input file which you are using in the code(dataset.csv) and also the screenshot of the complete error log.

Meanwhile, you can also refer to the following link where a similar issue is addressed

https://community.intel.com/t5/Intel-oneAPI-Data-Parallel-C/SYCL-kernel-cannot-call-an-undefined-function-without-SYCL/m-p/1382886#M2118

Regards,

Vidya.

[Jeremy_Wang]

Hi Vidya, 

Deeply appreciate for your help! 

Here is the full building log of the program. 

Build started...
1>------ Build started: Project: Macau_Entertainment_Index, Configuration: Debug x64 ------
1>In file included from Macau_Entertainment_Index.cpp:6:
1>In file included from E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl.hpp:11:
1>In file included from E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl/ONEAPI/atomic.hpp:11:
1>In file included from E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl/ONEAPI/atomic_accessor.hpp:11:
1>In file included from E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl/ONEAPI/atomic_enums.hpp:12:
1>In file included from E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl/access/access.hpp:10:
1>In file included from E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl/detail/common.hpp:122:
1>In file included from E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl/exception.hpp:16:
1>In file included from E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl/stl.hpp:16:
1>In file included from D:\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\include\functional:18:
1>In file included from D:\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\include\unordered_map:11:
1>In file included from D:\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\include\xhash:16:
1>D:\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\include\vector(1575,9): : error : SYCL kernel cannot call a variadic function
1>        _STL_VERIFY(
1>        ^
1>D:\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\include\yvals.h(178,13): note: expanded from macro '_STL_VERIFY'
1>            _STL_REPORT_ERROR(mesg);                                                       \
1>            ^
1>D:\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\include\yvals.h(169,9): note: expanded from macro '_STL_REPORT_ERROR'
1>        _RPTF0(_CRT_ASSERT, mesg);           \
1>        ^
1>C:\Program Files (x86)\Windows Kits\10\Include\10.0.19041.0\ucrt\crtdbg.h(764,37): note: expanded from macro '_RPTF0'
1>    #define _RPTF0(rptno, msg)      _RPT_BASE(rptno, __FILE__, __LINE__, NULL, "%s", msg)
1>                                    ^
1>C:\Program Files (x86)\Windows Kits\10\Include\10.0.19041.0\ucrt\crtdbg.h(751,48): note: expanded from macro '_RPT_BASE'
1>        (void) ((1 != _CrtDbgReport(__VA_ARGS__)) || \
1>                                               ^
1>Macau_Entertainment_Index.cpp(212,52): note: called by 'operator()'
1>                    ass_temp[index] = ass_data[row][col] * ass_weight[col];
1>                                                   ^
1>E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl/handler.hpp(2303,7): note: called by 'operator()'
1>      KernelFunc(Arg);
1>      ^
1>E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl/handler.hpp(899,5): note: called by 'kernel_parallel_for<sycl::detail::__pf_kernel_wrapper<(lambda at Macau_Entertainment_Index.cpp:209:17)>, sycl::id<2>, (lambda at E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl/handler.hpp:2299:12)>'
1>    KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()));
1>    ^
1>In file included from Macau_Entertainment_Index.cpp:6:
1>In file included from E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl.hpp:11:
1>In file included from E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl/ONEAPI/atomic.hpp:11:
1>In file included from E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl/ONEAPI/atomic_accessor.hpp:11:
1>In file included from E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl/ONEAPI/atomic_enums.hpp:12:
1>In file included from E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl/access/access.hpp:10:
1>In file included from E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl/detail/common.hpp:122:
1>In file included from E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl/exception.hpp:16:
1>In file included from E:\Intel\oneAPI\compiler\2021.3.0\windows\bin\..\include\sycl\CL/sycl/stl.hpp:16:
1>In file included from D:\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\include\functional:18:
1>In file included from D:\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\include\unordered_map:11:
1>In file included from D:\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\include\xhash:16:
1>D:\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\include\vector(1575,9): : error : SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute
1>        _STL_VERIFY(
1>        ^
1>D:\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\include\yvals.h(178,13): note: expanded from macro '_STL_VERIFY'
1>            _STL_REPORT_ERROR(mesg);                                                       \
1>            ^
1>D:\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\include\yvals.h(169,9): note: expanded from macro '_STL_REPORT_ERROR'
1>        _RPTF0(_CRT_ASSERT, mesg);           \
1>        ^
1>C:\Program Files (x86)\Windows Kits\10\Include\10.0.19041.0\ucrt\crtdbg.h(764,37): note: expanded from macro '_RPTF0'
1>    #define _RPTF0(rptno, msg)      _RPT_BASE(rptno, __FILE__, __LINE__, NULL, "%s", msg)
1>                                    ^
1>C:\Program Files (x86)\Windows Kits\10\Include\10.0.19041.0\ucrt\crtdbg.h(751,23): note: expanded from macro '_RPT_BASE'
1>        (void) ((1 != _CrtDbgReport(__VA_ARGS__)) || \
1>                      ^
1>C:\Program Files (x86)\Windows Kits\10\Include\10.0.19041.0\ucrt\crtdbg.h(611,26): note: '_CrtDbgReport' declared here
1>    _ACRTIMP int __cdecl _CrtDbgReport(
1>                         ^
1>D:\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\include\vector(1572,50): note: called by 'operator[]'
1>    _NODISCARD _CONSTEXPR20_CONTAINER const _Ty& operator[](const size_type _Pos) const noexcept /* strengthened */ {
1>                                                 ^
1>D:\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\include\vector(1575,9): : error : SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute
1>        _STL_VERIFY(
1>        ^
1>D:\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\include\yvals.h(178,13): note: expanded from macro '_STL_VERIFY'
1>            _STL_REPORT_ERROR(mesg);                                                       \
1>            ^
1>D:\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\include\yvals.h(170,9): note: expanded from macro '_STL_REPORT_ERROR'
1>        _CRT_SECURE_INVALID_PARAMETER(mesg); \
1>        ^
1>C:\Program Files (x86)\Windows Kits\10\Include\10.0.19041.0\ucrt\corecrt.h(371,15): note: expanded from macro '_CRT_SECURE_INVALID_PARAMETER'
1>            ::_invalid_parameter(_CRT_WIDE(#expr), __FUNCTIONW__, __FILEW__, __LINE__, 0)
1>              ^
1>C:\Program Files (x86)\Windows Kits\10\Include\10.0.19041.0\ucrt\corecrt.h(348,27): note: '_invalid_parameter' declared here
1>    _ACRTIMP void __cdecl _invalid_parameter(
1>                          ^
1>D:\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\include\vector(1572,50): note: called by 'operator[]'
1>    _NODISCARD _CONSTEXPR20_CONTAINER const _Ty& operator[](const size_type _Pos) const noexcept /* strengthened */ {
1>                                                 ^
1>3 errors generated.
1>Done building project "Macau_Entertainment_Index.vcxproj" -- FAILED.
========== Build: 0 succeeded, 1 failed, 0 up-to-date, 0 skipped ==========

About the environment, I'm currently using the Visual Studio 2019 community. And here is the version of DPC++ complier: 

Intel(R) oneAPI DPC++/C++ Compiler 2021.3.0 (2021.3.0.20210619)
Target: x86_64-pc-windows-msvc
Thread model: posix

The OS system I'm using is Windows 10 Pro for Workstations. The CPU is Intel(R) Core(TM) i7-7820HK CPU @ 2.90GHz 2.90 GHz. 

And here is the screen when I complie the program. The error E0751 sometimes will shows up. But as I clicked into the accessor_property_list.hpp, the error will disappear. 

And here is the code for the read() function. 

vector<vector<float>> read(string file) {
    vector<vector<float>>a;  //2-D vector storage in this variable
    vector<float>b;
    vector<string> row;
    string line;
    string filename;
    ifstream in(file);
    if (in.fail()) { cout << "File not found" << std::endl; return a; }
    while (getline(in, line))
    {
        file_to_string(row, line, ',');  //Take the int character in each line，"," as split sign
        for (int i = 0, leng = row.size(); i < leng; i++) {
            b.push_back(string_to_float(row[i]));
        }
        a.push_back(b);
        b.clear();
    }
    in.close();
    return a;
}

inline void file_to_string(vector<string>& record, const string& line, char delimiter) {
    int linepos = 0;
    char c;
    int linemax = line.length();
    string curstring;
    record.clear();
    while (linepos < linemax) {
        c = line[linepos];
        if (isdigit(c) || c == '.') {
            curstring += c;
        }
        else if (c == delimiter && curstring.size()) {
            record.push_back(curstring);
            curstring = "";
        }
        ++linepos;
    }
    if (curstring.size())
        record.push_back(curstring);
    return;
}

inline float string_to_float(string str) {
    int i = 0, len = str.length();
    float sum = 0;
    while (i < len) {
        if (str[i] == '.') break;
        sum = sum * 10 + str[i] - '0';
        ++i;
    }
    ++i;
    float t = 1, d = 1;
    while (i < len) {
        d *= 0.1;
        t = str[i] - '0';
        sum += t * d;
        ++i;
    }
    return sum;
}

And the dataset.csv file is in the attachment. 

Again, thank you for your kindly help! 

[VidyalathaB_Intel]

Hi Jeremy,

Thanks for sharing the input files and relevant functions.

Could you please update your dpcpp compiler to the latest version which is 2022.1 and check if you get the same error during compilation?

Please get back to us with the VS project file that you are working with if you encounter the same even with the latest compiler so that we could check it from our end as we are unable to reproduce the issue.

Additionally, could you please confirm if you are getting the same issue when trying to run it from Intel oneAPI command prompt?

Regards,

Vidya.

[Jeremy_Wang]

Hi Vidya, 

Sorry for late response. I reinstalled all the environment as you said. I upgraded the dpcpp into 2022.1 version and reinstall the VS2022 to see if it's the environment problem. The error still generate. And I tried to compile and run with the oneAPI command prompt, it seems work well in building process, but the calulation seems have something wrong. 

I'll uploadthe whole VS project file in the attachment. 

Thank you for your patience helping! 

[VidyalathaB_Intel]

Hi Jeremy,

Thanks for getting back to us.

>> but the calulation seems have something wrong. 

Could you please show us the results that you are getting and the expected output as well so that we could check it from our end?

Regards,

Vidya.

[Jeremy_Wang]

Hi Vidya, 

I double checked the result I mentioned. Seems like the it has no problem. And the program can run but might be stucked in some where. And I updated the source code file. Here should be the output. 

Weight for 0: 4.03226
Weight for 1: 5.50055
Weight for 2: 13.5135
Weight for 3: 5.9312
Weight for 4: 5.05561
Weight for 5: 6.41026

 I tried mannually compile it with the latest program. Building is okay. When running the output, seems like it stucked in somewhere. And when compile in Visual Studio, it still showing the same error when building. 

Thank you again for your patience! 

[VidyalathaB_Intel]

Hi Jeremy,

Thanks for the confirmation.

The output screen is getting hanged during the execution of the second q.submit statement. To debug the case, instead of the logic that you have written, I just loaded the ass_temp[index] variable with a constant value in order to understand the issue and I could see that it is working fine and there is no hanging this time.

In the second trial, this time I filled the ass_temp[index] variable with ass_weight[col] and this also worked.

So I suspect the issue comes with ass_data[row][col] this part of the logic.

Could you please cross-check the computation part of the code which is ass_temp[index] = ass_data[row][col] * ass_weight[col]; ?

And do let us know what you are trying to calculate from the second q.submit statement so that it will help us to address your issue accordingly.

I checked the output using the host_accessor as below (the same way how you have retrieved the values in the first q.submit)

host_accessor ass(buf_temp_MEI, read_only);

  for (int i = 0; i < data.size(); i++) {

    for (int j = 0; j < data[0].size(); j++)

    {

      cout << " ass_temp for 2nd submit " << i << ": "  << j << ": "<< ass[i][j] << std::endl;

    }

  }

Please do let me know if there is any misinterpretation of the scenario (feel free to correct it).

Regards,

Vidya.

[Jeremy_Wang]

Hi Vidya, 

Thank you for your kindly help. 

The vector data is storing the data of some stocks' price. Each column is a different stock. 

The second q.submit statement is mainly to calculate a weighted price according to the weight computed in the first q.submit. 

ass_temp[index] = ass_data[row][col] * ass_weight[col];
# This code can be also written as:
ass_temp[row][col] = ass_data[row][col] * ass_weight[col];

 The building error in VS is still exist. And I'm also quite sure that there are something wrong with the secomd submit(). 

Thank you for your time Vidya! 

Regards, 

Jeremy. 

[VidyalathaB_Intel]

Hi Jeremy,

>> The building error in VS is still exist

Ok, could you please try running your code in Release mode and see if you still observe the same error in VS?

Regards,

Vidya.

[VidyalathaB_Intel]

Hi Jeremy,

>>When running the output, seems like it stucked in somewhere. ...And I'm also quite sure that there are something wrong with the secomd submit(). 

We are working on your issue. we will get back to you soon.

Thanks for your patience.

Regards,

Vidya.

[Jeremy_Wang]

Hi Vidya, 

Thank you for your kindly help. It's quite a weriod problem. 

Looking forward to your reply. 

Thanks. 

Regards, 

Jeremy. 

[Alex_Y_Intel]

What building error did you see?

Are you intending to run the kernel on CPU or GPU?

[Jeremy_Wang]

Hi,

The building error I was already posted on the earlier reply to Vidya, don't know if you can see it. 

I was intending to run the kernel on FPGA emulator. And it should be run on CPU without real FPGA. 

[Alex_Y_Intel]

I need to clarify something here--while you're using "default_selector," how do you make sure the kernel is offloaded to CPU instead of GPU? The reason I asked was that if I purposely let it dispatch the kernel to CPU by changing the code to "cpu_selector," the result is very different from using "default_selector." Please to aware that when you use default_selector, in our case, it's defaulted to GPU, not CPU. Since you mentioned that you intend the kernel to run on FPGA, what's your expected answer on FPGA and what's the result now?

As for the build/compile issue, I can see the problem when I use "debug" mode, but "release" mode seems fine to me--I can build/compile the program successfully, is it the same as how you observed?

[Alex_Y_Intel]

Since there's no more update, I'm closing this from our site, but please feel free to discuss the issue with others in the forum or start another new topic if necessary.