<html xmlns:v="urn:schemas-microsoft-com:vml" xmlns:o="urn:schemas-microsoft-com:office:office" xmlns:w="urn:schemas-microsoft-com:office:word" xmlns:m="http://schemas.microsoft.com/office/2004/12/omml" xmlns="http://www.w3.org/TR/REC-html40"><head><META HTTP-EQUIV="Content-Type" CONTENT="text/html; charset=us-ascii"><meta name=Generator content="Microsoft Word 15 (filtered medium)"><style><!--
/* Font Definitions */
@font-face
{font-family:Calibri;
panose-1:2 15 5 2 2 2 4 3 2 4;}
/* Style Definitions */
p.MsoNormal, li.MsoNormal, div.MsoNormal
{margin:0cm;
margin-bottom:.0001pt;
font-size:11.0pt;
font-family:"Calibri","sans-serif";
mso-fareast-language:EN-US;}
a:link, span.MsoHyperlink
{mso-style-priority:99;
color:#0563C1;
text-decoration:underline;}
a:visited, span.MsoHyperlinkFollowed
{mso-style-priority:99;
color:#954F72;
text-decoration:underline;}
span.E-MailFormatvorlage17
{mso-style-type:personal-compose;
font-family:"Calibri","sans-serif";
color:windowtext;}
.MsoChpDefault
{mso-style-type:export-only;
font-family:"Calibri","sans-serif";
mso-fareast-language:EN-US;}
@page WordSection1
{size:612.0pt 792.0pt;
margin:70.85pt 70.85pt 2.0cm 70.85pt;}
div.WordSection1
{page:WordSection1;}
--></style><!--[if gte mso 9]><xml>
<o:shapedefaults v:ext="edit" spidmax="1026" />
</xml><![endif]--><!--[if gte mso 9]><xml>
<o:shapelayout v:ext="edit">
<o:idmap v:ext="edit" data="1" />
</o:shapelayout></xml><![endif]--></head><body lang=DE link="#0563C1" vlink="#954F72"><div class=WordSection1><p class=MsoNormal><span lang=EN-GB>Hello! <o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB><o:p> </o:p></span></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. It looks like the kernel template is not instantiated at all. Doing the instantiation by hand again works. <o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-GB>The used code: <o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-GB>#ifndef __CUDACC__<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-GB>#include <stddef.h><o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB>#define __constant__ __attribute__((constant))<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB>#define __device__ __attribute__((device))<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB>#define __global__ __attribute__((global))<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB>#define __host__ __attribute__((host))<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB>#define __shared__ __attribute__((shared))<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB>#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-GB>struct dim3 {<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB> unsigned x, y, z;<o:p></o:p></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) {}<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB>};<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-GB>typedef struct cudaStream *cudaStream_t;<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-GB>int cudaConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB> cudaStream_t stream = 0);<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-GB>#endif<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-GB><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-GB>template <typename T><o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB>__device__ int blubblub(T& a, float& b, double& c)<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB>{<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB> a = a * b;<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB> b = b - c;<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB> c = a * c;<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB> return a;<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB>}<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-GB><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-GB>template <typename T><o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB>__global__ void kernel(T a, float b, double c)<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB>{<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB> int result = blubblub<T>(a, b, c);<o:p></o:p></span></p><p class=MsoNormal>}<o:p></o:p></p><p class=MsoNormal><o:p> </o:p></p><p class=MsoNormal><o:p> </o:p></p><p class=MsoNormal>int main()<o:p></o:p></p><p class=MsoNormal>{<o:p></o:p></p><p class=MsoNormal><o:p> </o:p></p><p class=MsoNormal> kernel<int><<<dim3(1), dim3(1)>>>(5, 0.7f, 12.34);<o:p></o:p></p><p class=MsoNormal> <span lang=EN-GB>return 0;<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB>}<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-GB>The command line to compile: <o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB><o:p> </o:p></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 test.cu<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-GB>clang version 3.5 (trunk 200831)<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-GB><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-GB>Any help to fix this problem is highly appreciated. <o:p></o:p></span></p><p class=MsoNormal><span lang=EN-GB><o:p> </o:p></span></p><p class=MsoNormal><span lang=EN-GB>Michael Haidl<o:p></o:p></span></p><p class=MsoNormal><o:p> </o:p></p></div></body></html>