[LLVMdev] Generating PTX code from template kernel failes

[Justin Holewinski]

Looks like a Clang issue, adding cfe-dev.

I would file a bug report if you have not already.


On Thu, Feb 6, 2014 at 4:44 AM, Haidl, Michael <
michael.haidl at uni-muenster.de> wrote:

> 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.
>
>
>
> 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
>
> _______________________________________________
> LLVM Developers mailing list
> LLVMdev at cs.uiuc.edu         http://llvm.cs.uiuc.edu
> http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
>
>


-- 

Thanks,

Justin Holewinski
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20140211/609c5eeb/attachment.html>