Intel® oneAPI Base Toolkit
Support for the core tools and libraries within the base toolkit that are used to build and deploy high-performance data-centric applications.
418 Discussions

Problems with external const variables

eug
Beginner
1,310 Views
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.

0 Kudos
1 Solution
AbhishekD_Intel
Moderator
1,253 Views

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

 

View solution in original post

0 Kudos
5 Replies
AbhishekD_Intel
Moderator
1,290 Views

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



0 Kudos
eug
Beginner
1,277 Views

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:

Schermata 2020-10-27 alle 18.35.05.png

0 Kudos
AbhishekD_Intel
Moderator
1,254 Views

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

 

0 Kudos
eug
Beginner
1,242 Views
0 Kudos
AbhishekD_Intel
Moderator
1,232 Views

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


0 Kudos
Reply