[cfe-dev] Clang and CUDA with C++11 features
peter at colberg.org
Fri Jun 15 09:48:57 PDT 2012
On Fri, Jun 15, 2012 at 11:59:59AM -0400, Peter Colberg wrote:
> On Fri, Jun 15, 2012 at 08:31:50AM -0700, Manjunath Kudlur wrote:
> > As Justin mentioned, I don't think the right plumbing exists that uses the
> > frontend support for parsing and lowering the CUDA syntax and connects it
> > to the NVPTX backend. For instance, from the IR, it looks like the kernel
> > function pointer is just cast to i8* and passed to cudaLaunch. cudaLaunch
> > requires the name of the kernel to passed a char string.
> Definitely not, otherwise my cudaLaunch wrapper would be failing since 2007 ;-).
> The advantage of the CUDA runtime library, as opposed to the CUDA
> driver library, lies in the convenient execution of kernels.
> When the nvcc frontend converts the <<< >>> syntax to proper C++ code,
> it does not pass a char string to cudaLaunch, but an actual pointer.
> The only question is how to get that pointer in Clang…
I have to admit that I never looked at the memory pointed to by a
kernel function pointer, so I did (please tell me if too naïvely…).
__global__ void f(int* array)
array = 42;
void (*p)(int*) = &f;
printf("%s\n", reinterpret_cast<char const*>(p));
nvcc -Xcompiler -Wall -Xptxas -v -o function function.cu
ptxas info : Compiling entry function '_Z1fPi' for 'sm_10'
ptxas info : Used 2 registers, 8+16 bytes smem
So the function pointer to the GPU kernel is indeed not an ASCII string.
What is it then?
More information about the cfe-dev