Community
cancel
Showing results for 
Search instead for 
Did you mean: 
Highlighted
Beginner
165 Views

Problems with external const variables

Jump to solution
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

Accepted Solutions
Highlighted
Moderator
108 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
Highlighted
Moderator
145 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
Highlighted
Beginner
132 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
Highlighted
Moderator
109 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
Highlighted
Beginner
97 Views

Yes, it solves my issue.

Thank you.

0 Kudos
Highlighted
Moderator
87 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