<div dir="ltr">Looks like a Clang issue, adding cfe-dev.<div><br></div><div>I would file a bug report if you have not already.</div></div><div class="gmail_extra"><br><br><div class="gmail_quote">On Thu, Feb 6, 2014 at 4:44 AM, Haidl, Michael <span dir="ltr"><<a href="mailto:michael.haidl@uni-muenster.de" target="_blank">michael.haidl@uni-muenster.de</a>></span> wrote:<br>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"><div lang="DE" link="#0563C1" vlink="#954F72"><div><p class="MsoNormal">Hello! <u></u><u></u></p><p class="MsoNormal"><u></u> <u></u></p>
<p class="MsoNormal"><span lang="EN-GB">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.<u></u><u></u></span></p>
<p class="MsoNormal"><span lang="EN-GB"><u></u> <u></u></span></p><p class="MsoNormal"><span lang="EN-GB">The used code: <u></u><u></u></span></p><p class="MsoNormal"><span lang="EN-GB"><u></u> <u></u></span></p><p class="MsoNormal">
<span lang="EN-GB">#ifndef __CUDACC__<u></u><u></u></span></p><p class="MsoNormal"><span lang="EN-GB"><u></u> <u></u></span></p><p class="MsoNormal"><span lang="EN-GB">#include <stddef.h><u></u><u></u></span></p><p class="MsoNormal">
<span lang="EN-GB">#define __constant__ __attribute__((constant))<u></u><u></u></span></p><p class="MsoNormal"><span lang="EN-GB">#define __device__ __attribute__((device))<u></u><u></u></span></p><p class="MsoNormal"><span lang="EN-GB">#define __global__ __attribute__((global))<u></u><u></u></span></p>
<p class="MsoNormal"><span lang="EN-GB">#define __host__ __attribute__((host))<u></u><u></u></span></p><p class="MsoNormal"><span lang="EN-GB">#define __shared__ __attribute__((shared))<u></u><u></u></span></p><p class="MsoNormal">
<span lang="EN-GB">#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))<u></u><u></u></span></p><p class="MsoNormal"><span lang="EN-GB"><u></u> <u></u></span></p><p class="MsoNormal"><span lang="EN-GB">struct dim3 {<u></u><u></u></span></p>
<p class="MsoNormal"><span lang="EN-GB"> unsigned x, y, z;<u></u><u></u></span></p><p class="MsoNormal"><span lang="EN-GB"> __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {}<u></u><u></u></span></p>
<p class="MsoNormal"><span lang="EN-GB">};<u></u><u></u></span></p><p class="MsoNormal"><span lang="EN-GB"><u></u> <u></u></span></p><p class="MsoNormal"><span lang="EN-GB">typedef struct cudaStream *cudaStream_t;<u></u><u></u></span></p>
<p class="MsoNormal"><span lang="EN-GB"><u></u> <u></u></span></p><p class="MsoNormal"><span lang="EN-GB">int cudaConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,<u></u><u></u></span></p><p class="MsoNormal">
<span lang="EN-GB"> cudaStream_t stream = 0);<u></u><u></u></span></p><p class="MsoNormal"><span lang="EN-GB"><u></u> <u></u></span></p><p class="MsoNormal"><span lang="EN-GB">#endif<u></u><u></u></span></p>
<p class="MsoNormal"><span lang="EN-GB"><u></u> <u></u></span></p><p class="MsoNormal"><span lang="EN-GB"><u></u> <u></u></span></p><p class="MsoNormal"><span lang="EN-GB">template <typename T><u></u><u></u></span></p>
<p class="MsoNormal"><span lang="EN-GB">__device__ int blubblub(T& a, float& b, double& c)<u></u><u></u></span></p><p class="MsoNormal"><span lang="EN-GB">{<u></u><u></u></span></p><p class="MsoNormal"><span lang="EN-GB"> a = a * b;<u></u><u></u></span></p>
<p class="MsoNormal"><span lang="EN-GB"> b = b - c;<u></u><u></u></span></p><p class="MsoNormal"><span lang="EN-GB"> c = a * c;<u></u><u></u></span></p><p class="MsoNormal"><span lang="EN-GB"> return a;<u></u><u></u></span></p>
<p class="MsoNormal"><span lang="EN-GB">}<u></u><u></u></span></p><p class="MsoNormal"><span lang="EN-GB"><u></u> <u></u></span></p><p class="MsoNormal"><span lang="EN-GB"><u></u> <u></u></span></p><p class="MsoNormal"><span lang="EN-GB">template <typename T><u></u><u></u></span></p>
<p class="MsoNormal"><span lang="EN-GB">__global__ void kernel(T a, float b, double c)<u></u><u></u></span></p><p class="MsoNormal"><span lang="EN-GB">{<u></u><u></u></span></p><p class="MsoNormal"><span lang="EN-GB"> int result = blubblub<T>(a, b, c);<u></u><u></u></span></p>
<p class="MsoNormal">}<u></u><u></u></p><p class="MsoNormal"><u></u> <u></u></p><p class="MsoNormal"><u></u> <u></u></p><p class="MsoNormal">int main()<u></u><u></u></p><p class="MsoNormal">{<u></u><u></u></p><p class="MsoNormal">
<u></u> <u></u></p><p class="MsoNormal"> kernel<int><<<dim3(1), dim3(1)>>>(5, 0.7f, 12.34);<u></u><u></u></p><p class="MsoNormal"> <span lang="EN-GB">return 0;<u></u><u></u></span></p>
<p class="MsoNormal"><span lang="EN-GB">}<u></u><u></u></span></p><p class="MsoNormal"><span lang="EN-GB"><u></u> <u></u></span></p><p class="MsoNormal"><span lang="EN-GB">The command line to compile: <u></u><u></u></span></p>
<p class="MsoNormal"><span lang="EN-GB"><u></u> <u></u></span></p><p class="MsoNormal"><span lang="EN-GB">clang++ -x cuda -S -emit-llvm -target nvptx64 -Xclang -fcuda-is-device -o test.dev.ll <a href="http://test.cu" target="_blank">test.cu</a><u></u><u></u></span></p>
<p class="MsoNormal"><span lang="EN-GB"><u></u> <u></u></span></p><p class="MsoNormal">clang version 3.5 (trunk 200831)<u></u><u></u></p><p class="MsoNormal"><u></u> <u></u></p><p class="MsoNormal"><u></u> <u></u></p><p class="MsoNormal">
<span lang="EN-GB">Any help to fix this problem is highly appreciated. <span class="HOEnZb"><font color="#888888"><u></u><u></u></font></span></span></p><span class="HOEnZb"><font color="#888888"><p class="MsoNormal"><span lang="EN-GB"><u></u> <u></u></span></p>
<p class="MsoNormal"><span lang="EN-GB">Michael Haidl<u></u><u></u></span></p></font></span></div></div><br>_______________________________________________<br>
LLVM Developers mailing list<br>
<a href="mailto:LLVMdev@cs.uiuc.edu">LLVMdev@cs.uiuc.edu</a> <a href="http://llvm.cs.uiuc.edu" target="_blank">http://llvm.cs.uiuc.edu</a><br>
<a href="http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev" target="_blank">http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev</a><br>
<br></blockquote></div><br><br clear="all"><div><br></div>-- <br><br><div>Thanks,</div><div><br></div><div>Justin Holewinski</div>
</div>