<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:"Cambria Math";
panose-1:2 4 5 3 5 4 6 3 2 4;}
@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>Hello! <o:p></o:p></p><p class=MsoNormal><o:p> </o:p></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.<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>clang version 3.5 (trunk 200831)<o:p></o:p></p><p class=MsoNormal><o:p> </o:p></p><p class=MsoNormal><o:p> </o:p></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></div></body></html>