[LLVMbugs] [Bug 18778] New: clang fails to emit LLVM for template CUDA kernels (device code with -fcuda-is-device)

[bugzilla-daemon at llvm.org]

http://llvm.org/bugs/show_bug.cgi?id=18778

            Bug ID: 18778
           Summary: clang fails to emit LLVM for template CUDA kernels
                    (device code with -fcuda-is-device)
           Product: clang
           Version: trunk
          Hardware: All
                OS: All
            Status: NEW
          Severity: normal
          Priority: P
         Component: LLVM Codegen
          Assignee: unassignedclangbugs at nondot.org
          Reporter: michael.haidl at uni-muenster.de
                CC: llvmbugs at cs.uiuc.edu
    Classification: Unclassified

When compiling valid CUDA code to LLVM (clang++ -Xclang -fcuda-is-device -S
-emit-llvm -target nvptx64) clang fails to emit LLVM code on template kernels.
This may happen because the global attribute get lost on template
instanciation. Non-template kernels compile fine and working LLVM code is
emitted. However, template device functions called from a global function is
correctly instanciated and emitted.

Code to reproduce:

#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))

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 dev_f(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 = dev_f<T>(a, b, c);
}

int main()
{
        kernel<int><<<dim3(1), dim3(1)>>>(5, 0.7f, 12.34);
        return 0;
}

-- 
You are receiving this mail because:
You are on the CC list for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-bugs/attachments/20140208/494b3f2c/attachment.html>