Intel® oneAPI Data Parallel C++
Support for Intel® oneAPI DPC++ Compiler, Intel® oneAPI DPC++ Library, Intel ICX Compiler , Intel® DPC++ Compatibility Tool, and GDB*

Calling a function in kernel scope?

Ajmeer_Khan
Employee
1,969 Views

Hi ,

I am new to dpc++ program. I have attached my code below. I want to call user function (add) from kernel function (parallel_for) with an accessor as a parameter.

 

Since the argument in the user function is a pointer I am unable to map the accessor to it.

Can someone get me a way out.

Thanks,

Ajmeer Khan.

0 Kudos
1 Solution
Ajmeer_Khan
Employee
1,928 Views

Thanks Goutham . 

It worked perfectly.  I have few doubts.

1. How can I access the global variable(data) from kernel function?

2.How to return the value(from add function) to kernel function?

View solution in original post

0 Kudos
4 Replies
GouthamK_Intel
Moderator
1,958 Views

Hi Ajmeer Khan,

We saw the attached source code, you are calling the add function inside the kernel by passing accessor as a parameter, but on the other side in the function definition, it is expecting an int* type value. Hence you are getting errors.

You can overcome this problem in two ways.

1) Use accessor datatype instead of int* variable inside the function definition /declaration.

Example:

void add(cl::sycl::accessor<int, 1, cl::sycl::access::mode::write> p, int x){
    p[x]=10;
}

 

2) Typecast accessor address (aa) to (int *) inside the kernel and use int * in add function definition/ declaration.

Example:

//Kernel
h.parallel_for(range<1>{10},[=](id<1> item)
      {
        add((int *)&aa[0],3);
      });

//Outside Main()
void add(int *p,int x)
{
  p[x]=10;
}

 

Also, I'm attaching the complete working code sample below.

#include<iostream>
#include<CL/sycl.hpp>

using namespace std;
using namespace sycl;

//void add(int *p, int x);
void add(cl::sycl::accessor<int, 1, cl::sycl::access::mode::write> p, int x);

int main()
{
    int a[10],b[10],c[10];
    for(int i=0;i<10;i++)
    {
        a[i]=i+1;
        b[i]=2;
    }
    cpu_selector device_selector;
    queue q(device_selector);
    {
        buffer<int ,1> a_buff(a,range<1> (10));
       // buffer<int ,1> b_buff(b,range<1> (10));
        //buffer<int,1> c_buff(c,range<1> (10));
        cout<<q.get_device().get_info<info::device::name>().c_str() << "\n";
        q.submit([&](handler &h)
        {
            cl::sycl::stream kernelout(512*1024, 512, h);
            auto aa=a_buff.get_access<access::mode::write>(h);
           // auto bb=b_buff.get_access<access::mode::write>(h);
            //auto cc=c_buff.get_access<access::mode::write>(h);
            h.parallel_for(range<1>{10},[=](id<1> item)
            {
                add(aa,3);
                 //add((int *)&aa[0],3);
            });
      });
    }
    for(int i=0;i<10;i++)
    {
        cout<<a[i]<<"\t";
    }
}

void add(cl::sycl::accessor<int, 1, cl::sycl::access::mode::write> p, int x){
        p[x]=10;
}

/*
void add(int *p,int x)
{
    p[x]=10;
}
*/

 

Note: We observed that you are using SYCL_EXTERNAL, but in your specific code you don't need to mark your function as SYCL_EXTERNAL.

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

You can refer the below link for more information.

https://community.intel.com/t5/Intel-oneAPI-Data-Parallel-C/Calling-a-function-in-kernel-scope/td-p/1184545

 

Please feel free to reach out to us if you are facing any further issues in running this code.

 

Thanks & Regards

Goutham

 

0 Kudos
Ajmeer_Khan
Employee
1,929 Views

Thanks Goutham . 

It worked perfectly.  I have few doubts.

1. How can I access the global variable(data) from kernel function?

2.How to return the value(from add function) to kernel function?

0 Kudos
GouthamK_Intel
Moderator
1,921 Views

Hi Ajmeer Khan,

Thanks for the confirmation.

Glad to know that your issue is resolved!


As the primary issue for which you had raised this thread is resolved, we would appreciate if you can raise a new thread for any other issues/doubts you have and we will be very much happy to resolve your issues. Also, it would be easy for other community members and us to keep track of the threads.

Have a Good day!


Thanks & Regards

Goutham


0 Kudos
GouthamK_Intel
Moderator
1,901 Views

Hi Ajmeer Khan,

As this issue has been resolved, we will no longer respond to this thread. 

If you require any additional assistance from Intel, please start a new thread. 

Any further interaction in this thread will be considered community only. 

Have a Good day!


Thanks & Regards

Goutham


0 Kudos
Reply