PGI User Forum
 SearchSearch   MemberlistMemberlist     RegisterRegister   ProfileProfile    Log inLog in 

Free OpenACC Webinar

Is it possible to call a CUDA kernel from PGI compiled code?
Goto page 1, 2  Next
 
Post new topic   Reply to topic    PGI User Forum Forum Index -> Accelerator Programming
View previous topic :: View next topic  
Author Message
PaulPa



Joined: 02 Aug 2012
Posts: 35

PostPosted: Sun Nov 25, 2012 6:57 am    Post subject: Is it possible to call a CUDA kernel from PGI compiled code? Reply with quote

Hi,

I'm just curious if the statement of PGI's OpenACC FAQ ist still valid:

"PGI is working on the design of a feature to allow you to call kernel functions written in CUDA or PTX or other languages directly from your C or Fortran program. We will announce this feature when it is available."

If this is still up-to-date, could you please tell me when this feature is likely
to be added?

I'm looking forward to this feature since it would allow me to "hand tune"
some "hot" kernels.

Best,
Paul
Back to top
View user's profile
mkcolg



Joined: 30 Jun 2004
Posts: 6213
Location: The Portland Group Inc.

PostPosted: Mon Nov 26, 2012 12:34 pm    Post subject: Reply with quote

Hi Paul,

This is part of the proposed OpenACC 2.0 Spec. See: http://www.openacc.org/node/173 and http://www.openacc.org/sites/default/files/Proposed%20Additions%20for%20OpenACC%202.pdf. In particular, see the "routine" directive.

We expect to have these features implemented by mid next year.

- Mat
Back to top
View user's profile
PaulPa



Joined: 02 Aug 2012
Posts: 35

PostPosted: Mon Nov 26, 2012 12:54 pm    Post subject: Reply with quote

Hi Mat,

as always, thank you :)

EDIT: A closer look into the OpenACC_2.pdf reveals that we might be talking about two different things. You are talking about routines that can be called within an open acc region, right?

I was talking about a cuda kernel that can be called outside of an acc region.
E.g.:
Code:

some_cuda_kernel<<<num_blocks, threads_per_block>>>(in_out);
    // do something before acc region
#pragma acc parallel deviceptr(in_out)
    //do something with in_out
//copy in_out to C code


Best,
Paul
Back to top
View user's profile
mkcolg



Joined: 30 Jun 2004
Posts: 6213
Location: The Portland Group Inc.

PostPosted: Mon Nov 26, 2012 3:10 pm    Post subject: Reply with quote

Hi Paul,

Quote:
You are talking about routines that can be called within an open acc region, right?
Correct. I'll see if we can make the FAQ more clear that it's about calling CUDA device functions from within an OpenACC compute region.

Quote:
I was talking about a cuda kernel that can be called outside of an acc region.
With Fortran you can do this now by using PGI CUDA Fortran.

As for C, the problem has to do with NVIDIA's header files where they need to make some changes to allow PGCC to be used as a host compiler. We've asked several times, but they haven't as of yet. Feel free to put a request in to NVIDIA. Not sure it will help, but wouldn't hurt.

In the mean time, what you need to do is have your PGCC compiled code call your NVCC compiled CUDA C code.

- Mat
Back to top
View user's profile
PaulPa



Joined: 02 Aug 2012
Posts: 35

PostPosted: Tue Nov 27, 2012 2:46 pm    Post subject: Reply with quote

My project is structured as follows:
- main.c
- compute.c (this file issues calls to the cuda runtime API)
- kernels.cu

I can compile this project with Intel's compiler for the *.c files and NVIDIA's nvcc for the *.cu file and finally linking everything with Intel's compiler.

If I try exactly the same with the pgi compiler I receive the following error, while compiling the compute.c file:
Code:
PGC-F-0249-#error --  --- !!! UNKNOWN COMPILER: please provide a CUDA compatible definition for '__align__' !!! --- (/usr/local/cuda/5.0.35/include/host_defines.h: 128)


That's what you are talking about, right?

Quote:
In the mean time, what you need to do is have your PGCC compiled code call your NVCC compiled CUDA C code.


Please correct me if I'm mistaken, my compute.c file is not allowed to include the cuda_runtime.h and make cuda api calls? So I have to move all
cuda api calls to a different file which needs to be compiled with the nvcc?

Best,
Paul
Back to top
View user's profile
Display posts from previous:   
Post new topic   Reply to topic    PGI User Forum Forum Index -> Accelerator Programming All times are GMT - 7 Hours
Goto page 1, 2  Next
Page 1 of 2

 
Jump to:  
You cannot post new topics in this forum
You cannot reply to topics in this forum
You cannot edit your posts in this forum
You cannot delete your posts in this forum
You cannot vote in polls in this forum


Powered by phpBB © phpBB Group