Community
cancel
Showing results for 
Search instead for 
Did you mean: 
Highlighted
Employee
166 Views

Calling a function in kernel scope?

Jump to solution

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

Accepted Solutions
Highlighted
Employee
125 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
Highlighted
Moderator
155 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/...

 

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

 

Thanks & Regards

Goutham

 

0 Kudos
Highlighted
Employee
126 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
Highlighted
Moderator
118 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
Highlighted
Moderator
98 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