[cfe-dev] Clang/LLVM doesn't emit IR for cuda template kernel

[Haidl, Michael]

Hello! 

 

I'm trying to generate IR from CUDA C++. This works fine until templates
come into play. Using a  template __device__ function within a __global__
function works well. The specific instantiations of the template function
are generated. However, trying to forward a template parameter from the
kernel launch code to the device function breaks somehow the transformation
process and an empty .ll file is emitted. It looks like the kernel template
is not instantiated at all. Doing the instantiation by hand again works. 

 

The used code: 

 

#ifndef __CUDACC__

 

#include <stddef.h>

#define __constant__ __attribute__((constant))

#define __device__ __attribute__((device))

#define __global__ __attribute__((global))

#define __host__ __attribute__((host))

#define __shared__ __attribute__((shared))

#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))

 

struct dim3 {

  unsigned x, y, z;

  __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) :
x(x), y(y), z(z) {}

};

 

typedef struct cudaStream *cudaStream_t;

 

int cudaConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,

                      cudaStream_t stream = 0);

 

#endif

 

 

template <typename T>

__device__ int blubblub(T& a, float& b, double& c)

{

        a = a * b;

        b = b - c;

        c = a * c;

        return a;

}

 

 

template <typename T>

__global__ void kernel(T a, float b, double c)

{

        int result = blubblub<T>(a, b, c);

}

 

 

int main()

{

 

        kernel<int><<<dim3(1), dim3(1)>>>(5, 0.7f, 12.34);

        return 0;

}

 

The command line to compile: 

 

clang++ -x cuda -S -emit-llvm -target nvptx64 -Xclang -fcuda-is-device -o
test.dev.ll test.cu

 

clang version 3.5 (trunk 200831)

 

 

Any help to fix this problem is highly appreciated. 

 

Michael Haidl

 

-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-dev/attachments/20140207/c9ed5360/attachment.html>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: smime.p7s
Type: application/pkcs7-signature
Size: 6035 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-dev/attachments/20140207/c9ed5360/attachment.bin>