[cfe-dev] Clang and CUDA with C++11 features

[Peter Colberg]

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…).

// function.cu

__global__ void f(int* array)
{
    array[0] = 42;
}

int main()
{
    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


./function
UH��H��H�}�H�E�H�������UH��SH��(�P@

So the function pointer to the GPU kernel is indeed not an ASCII string.

What is it then?

Peter