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: [email protected]
Reporter: [email protected]
CC: [email protected]
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.
_______________________________________________
LLVMbugs mailing list
[email protected]
http://lists.cs.uiuc.edu/mailman/listinfo/llvmbugs
