Problems with external const variables

[eug]

Hello,

I have some problems to make a dpcpp kernel working if it uses external const variables defined in another cpp file.

For example, I have the following files:

// constants.cpp file:

#include "constants.hpp"

const int N=10;

// constant.hpp file:

extern const int N;

//main.cpp

#include "constant.hpp"

...

host code

device code

...

The problem is that the host program is able to find these variables but the kernel throw an exception like: 

" terminate called after throwing an instance of 'cl::sycl::compile_program_error'

 what():  The program was built for 1 devices

Build program log for 'Intel(R) Gen9 HD Graphics NEO':

error: undefined reference to `N'

error: backend compiler failed build.

-17 (CL_LINK_PROGRAM_FAILURE)

Aborted"

Thank you.

[AbhishekD_Intel]

Hi,

Please provide us a small sample of kernel code so that we will get the idea of where and how you are using this const variable. This will help us to give you details on why it's working with the host and not with iGPU.

It might be the case that you are directly passing the reference of this variable to the kernel so it's not supporting on the iGPU side.

So please send us a small reproducer of this issue.

Warm Regards,

Abhishek

[eug]

The sample includes 3 files: constants.hpp, constants.cpp and main.cpp. As you can see in the main file the constant MAX_N is printed correctly; instead the kernel throws an exception.

#ifndef CONSTANTS_HPP
#define CONSTANTS_HPP
#include <cstddef>
namespace constants {
    extern const size_t MAX_N;
};
#endif 

#include "constants.hpp"
namespace constants{
    const size_t MAX_N=7;
}

#include <CL/sycl.hpp>
#include "constants.hpp"

using namespace std;
using namespace cl::sycl;
using namespace constants;

void my_kernel(queue q, vector<int> &input){

    buffer<int> buff(input.data(), input.size());
    q.submit([&](sycl::handler& cgh) {

        auto acc_set=buff.get_access<access::mode::write>(cgh);
        cgh.parallel_for<class partition_kernel>(cl::sycl::range<1>{input.size()},[=](cl::sycl::id<1> idx) {
                        
            int sum=0;
            for(int i=0; i<MAX_N; i++){
                sum+=i;
            }
            acc_set[idx[0]] = sum;
        });
    }).wait();
}

int main(int argc, char **argv) {

	auto asyncHandler = [&](cl::sycl::exception_list eL) {
		for (auto& e : eL) {
			try {
				std::rethrow_exception(e);
			}catch (cl::sycl::exception& e) {
				std::cout << e.what() << std::endl;
				std::cout << "fail" << std::endl;
				// std::terminate() will exit the process, return non-zero, and output a
				// message to the user about the exception
				std::terminate();
			}
		}
	};

    cl::sycl::queue q(gpu_selector{});

    std::cout << "Runnning on " << q.get_device().get_info<sycl::info::device::name>() << "\n";
    
    cout<<"MAX_N value: "<<MAX_N<<std::endl;
    vector<int> input;
    
    for(int i=0; i<MAX_N; i++){
        input.push_back(0);
    }
    
    for(int j=0; j<MAX_N; j++){
        std::cout<<input[j]<<" ";
    }
    cout<<std::endl;

    my_kernel(q,input);
    
    for(int i=0; i<MAX_N; i++){
      std::cout<<input[i]<<" ";
    }     
    cout<<std::endl;
    return 0;
}

That's the exception:

[AbhishekD_Intel · Accepted solution]

Hi,

From your provided codes I can see that you have created the extern variable MAX_N, so its storage space will be the static memory from the host side. But you are trying to access the same variable inside the kernel which will run on iGPU side.

So in such a case as a workaround, you can capture that variable into some other variable that is known to iGPU. You can try the following way to run the same code on iGPU.

#include <CL/sycl.hpp>
#include "constants.hpp"

using namespace std;
using namespace cl::sycl;
using namespace constants;

void my_kernel(queue q, vector<int> &input){

    buffer<int> buff(input.data(), input.size());
    q.submit([&](sycl::handler& cgh) {

        auto acc_set=buff.get_access<access::mode::write>(cgh);
        cgh.parallel_for<class partition_kernel>(cl::sycl::range<1>{input.size()},[=, MAX_N_local=MAX_N](cl::sycl::id<1> idx) {

            int sum=0;
            for(int i=0; i<MAX_N_local; i++){
                sum+=i;
            }
            acc_set[idx[0]] = sum;
        });
    }).wait();
}

int main(int argc, char **argv) {

        auto asyncHandler = [&](cl::sycl::exception_list eL) {
                for (auto& e : eL) {
                        try {
                                std::rethrow_exception(e);
                        }catch (cl::sycl::exception& e) {
                                std::cout << e.what() << std::endl;
                                std::cout << "fail" << std::endl;
                                // std::terminate() will exit the process, return non-zero, and output a
                                // message to the user about the exception
                                std::terminate();
                        }
                }
        };

    cl::sycl::queue q(gpu_selector{});

    std::cout << "Runnning on " << q.get_device().get_info<sycl::info::device::name>() << "\n";

    cout<<"MAX_N value: "<<MAX_N<<std::endl;
    vector<int> input;

    for(int i=0; i<MAX_N; i++){
        input.push_back(0);
    }

    for(int j=0; j<MAX_N; j++){
        std::cout<<input[j]<<" ";
    }
    cout<<std::endl;

    my_kernel(q,input);

    for(int i=0; i<MAX_N; i++){
      std::cout<<input[i]<<" ";
    }
    cout<<std::endl;
    return 0;
}

I hope, the provided information helped you in solving your issue. Let us know if you have any issues related to it.

Warm Regards,

Abhishek

[eug]

Yes, it solves my issue.

Thank you.

[AbhishekD_Intel]

Hi,

Thanks for the confirmation. We will no longer monitor this thread, please raise a new thread if you have any issues.

Warm Regards,

Abhishek