<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=utf-8">
<meta name="Generator" content="Microsoft Word 14 (filtered medium)">
<style><!--
/* Font Definitions */
@font-face
{font-family:Calibri;
panose-1:2 15 5 2 2 2 4 3 2 4;}
@font-face
{font-family:Tahoma;
panose-1:2 11 6 4 3 5 4 4 2 4;}
@font-face
{font-family:Consolas;
panose-1:2 11 6 9 2 2 4 3 2 4;}
/* Style Definitions */
p.MsoNormal, li.MsoNormal, div.MsoNormal
{margin:0cm;
margin-bottom:.0001pt;
font-size:12.0pt;
font-family:"Times New Roman","serif";}
a:link, span.MsoHyperlink
{mso-style-priority:99;
color:blue;
text-decoration:underline;}
a:visited, span.MsoHyperlinkFollowed
{mso-style-priority:99;
color:purple;
text-decoration:underline;}
span.E-MailFormatvorlage17
{mso-style-type:personal-reply;
font-family:"Calibri","sans-serif";
color:#1F497D;}
.MsoChpDefault
{mso-style-type:export-only;
font-family:"Calibri","sans-serif";}
@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="EN-US" link="blue" vlink="purple">
<div class="WordSection1">
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1F497D">Hi Simon,<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1F497D"><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1F497D">I made a workaround by including a explicit copying from device to host after the kernel is called. The function looks now like this:<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1F497D"><o:p> </o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:blue">void</span><span style="font-size:9.5pt;font-family:Consolas;color:black"><o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">CUDA_gradient</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[3],<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">spacing</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[3],<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dev_in</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dev_out_x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dev_out_y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dev_out_z</span><span style="font-size:9.5pt;font-family:Consolas;color:black">)<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">{<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">int3</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> dev_Size
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">=</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">make_int3</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[0],
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[1],
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[2]);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">float3</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> dev_Spacing
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">=</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">make_float3</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">spacing</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[0],
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">spacing</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[1],
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">spacing</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[2]);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><o:p> </o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:green">// Output volume</span><span style="font-size:9.5pt;font-family:Consolas;color:black"><o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">long</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> outputMemorySize =
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[0] *
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[1] *
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[2] *
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">sizeof</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black">);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:black"> * dev_grad_x;<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:black"> * dev_grad_y;<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> * dev_grad_z;<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaMalloc</span><span style="font-size:9.5pt;font-family:Consolas;color:black">((</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">void</span><span style="font-size:9.5pt;font-family:Consolas;color:black">**)&dev_grad_x,
outputMemorySize);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaMalloc</span><span style="font-size:9.5pt;font-family:Consolas;color:black">((</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">void</span><span style="font-size:9.5pt;font-family:Consolas;color:black">**)&dev_grad_y,
outputMemorySize);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaMalloc</span><span style="font-size:9.5pt;font-family:Consolas;color:black">((</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">void</span><span style="font-size:9.5pt;font-family:Consolas;color:black">**)&dev_grad_z,
outputMemorySize);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaMemset</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(dev_grad_x, 2.f, outputMemorySize);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaMemset</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(dev_grad_y, 3.f, outputMemorySize);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaMemset</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(dev_grad_z, 4.f, outputMemorySize);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><o:p> </o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:green">// Thread Block Dimensions</span><span style="font-size:9.5pt;font-family:Consolas;color:black"><o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">dim3</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> dimBlock
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">=</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">dim3(</span><span style="font-size:9.5pt;font-family:Consolas;color:black">10, 10, 10</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">)</span><span style="font-size:9.5pt;font-family:Consolas;color:black">;<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><o:p> </o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> blocksInX =
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">iDivUp</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[0],
dimBlock.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> blocksInY =
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">iDivUp</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[1],
dimBlock.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> blocksInZ =
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">iDivUp</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[2],
dimBlock.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">z</span><span style="font-size:9.5pt;font-family:Consolas;color:black">);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><o:p> </o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">dim3</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> dimGrid
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">=</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">dim3(</span><span style="font-size:9.5pt;font-family:Consolas;color:black">blocksInX, blocksInY, blocksInZ</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">)</span><span style="font-size:9.5pt;font-family:Consolas;color:black">;<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><o:p> </o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">gradient_kernel</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> <<< dimGrid, dimBlock >>> (dev_in, dev_grad_x, dev_grad_y, dev_grad_z, dev_Size, dev_Spacing);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:#6F008A">CUDA_CHECK_ERROR</span><span style="font-size:9.5pt;font-family:Consolas;color:black">;<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><o:p> </o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaMemcpy</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dev_out_x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
dev_grad_x, outputMemorySize, </span><span style="font-size:9.5pt;font-family:Consolas;color:#314F4F">cudaMemcpyDeviceToHost</span><span style="font-size:9.5pt;font-family:Consolas;color:black">);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:#6F008A">CUDA_CHECK_ERROR</span><span style="font-size:9.5pt;font-family:Consolas;color:black">;<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaMemcpy</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dev_out_y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
dev_grad_y, outputMemorySize, </span><span style="font-size:9.5pt;font-family:Consolas;color:#314F4F">cudaMemcpyDeviceToHost</span><span style="font-size:9.5pt;font-family:Consolas;color:black">);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:#6F008A">CUDA_CHECK_ERROR</span><span style="font-size:9.5pt;font-family:Consolas;color:black">;<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaMemcpy</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dev_out_z</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
dev_grad_z, outputMemorySize, </span><span style="font-size:9.5pt;font-family:Consolas;color:#314F4F">cudaMemcpyDeviceToHost</span><span style="font-size:9.5pt;font-family:Consolas;color:black">);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:#6F008A">CUDA_CHECK_ERROR</span><span style="font-size:9.5pt;font-family:Consolas;color:black">;<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:9.5pt;font-family:Consolas;color:black">}<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:9.5pt;font-family:Consolas;color:black"><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:black">I am not sure if this violates the intended behavior of the ITKCudaCommon by explicitly copying the memory (that needed to be allocated before).<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:black"><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:black">Still I cannot solve why the implicit memory copying mechanism do not work.<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:black">Have you look into this problem or do you miss any information?<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:black">With best regards,<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:black">Gordian</span><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1F497D"><o:p></o:p></span></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal"><b><span lang="DE" style="font-size:10.0pt;font-family:"Tahoma","sans-serif"">Von:</span></b><span lang="DE" style="font-size:10.0pt;font-family:"Tahoma","sans-serif""> Simon Rit [mailto:simon.rit@creatis.insa-lyon.fr]
<br>
<b>Gesendet:</b> Mittwoch, 27. Februar 2019 21:58<br>
<b>An:</b> Kabelitz, Gordian<br>
<b>Cc:</b> rtk-users@public.kitware.com<br>
<b>Betreff:</b> Re: [Rtk-users] GPU kernel do not change output variables<o:p></o:p></span></p>
<p class="MsoNormal"><o:p> </o:p></p>
<div>
<div>
<p class="MsoNormal">Hi,<o:p></o:p></p>
</div>
<div>
<p class="MsoNormal">Sounds like a challenge. When you say you set fixed numbers, did you check that you reach the point where you set this number? You can use cuprintf to check what's going on in the kernel.<o:p></o:p></p>
</div>
<div>
<p class="MsoNormal">One thing wrong I noticed: you use size.Fill in a loop, which is a bit odd because it will Fill the size with the last value of the loop.<o:p></o:p></p>
</div>
<div>
<p class="MsoNormal">I hope this helps,<o:p></o:p></p>
</div>
<div>
<p class="MsoNormal">Simon<o:p></o:p></p>
</div>
</div>
<p class="MsoNormal"><o:p> </o:p></p>
<div>
<div>
<p class="MsoNormal">On Wed, Feb 27, 2019 at 9:39 PM Kabelitz, Gordian <<a href="mailto:Gordian.Kabelitz@medma.uni-heidelberg.de">Gordian.Kabelitz@medma.uni-heidelberg.de</a>> wrote:<o:p></o:p></p>
</div>
<blockquote style="border:none;border-left:solid #CCCCCC 1.0pt;padding:0cm 0cm 0cm 6.0pt;margin-left:4.8pt;margin-right:0cm">
<div>
<div>
<p class="MsoNormal" style="mso-margin-top-alt:auto;margin-bottom:3.75pt;line-height:12.0pt">
Hi rtk-users,<o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;margin-bottom:3.75pt;line-height:12.0pt">
I am facing an oddity which I cannot explain.<o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;margin-bottom:3.75pt;line-height:12.0pt">
I want to implement a new gradient filter. The input is an CudaImage<float,3> and the output should be an CudaImage<CovariantVector<float,3>,3>. The filter runs without any cuda errors but the output (pout_(xyz)) is has not changed at all. The kernel function
is accessed and the print out from there seems to be okay. I tried to explicitly copy the content of the GPUBuffer into the CPUBuffer. Still no success. Even if I set fixed numbers in the kernel to the output image nothing changed. I use CUDA 9.0, Visual
Studio 2015, ITK 5.0, RTK 2.0 as remote module, CMake 3.13., Windows 7 64bit. The relevant code snippets are below.<o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;margin-bottom:3.75pt;line-height:12.0pt">
Do I miss something obvious? Any recommendation are welcome.<o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;margin-bottom:3.75pt;line-height:12.0pt">
With kind regards,<o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;margin-bottom:3.75pt;line-height:12.0pt">
Gordian<o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"> <o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto">The GPUGenerateData function:<o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">GPUGenerateData()</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> {</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> inputSize[3];</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> outputSize[3];</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> inputSpacing[3];</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> outputSpacing[3];</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">for</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> (</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
i = 0; i<3; i++)</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> {</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> inputSize[i] =
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">this</span><span style="font-size:9.5pt;font-family:Consolas;color:black">->GetInput()->GetBufferedRegion().GetSize()[i];</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> outputSize[i] =
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">this</span><span style="font-size:9.5pt;font-family:Consolas;color:black">->GetOutput()->GetBufferedRegion().GetSize()[i];</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> inputSpacing[i] =
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">this</span><span style="font-size:9.5pt;font-family:Consolas;color:black">->GetInput()->GetSpacing()[i];</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> outputSpacing[i] =
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">this</span><span style="font-size:9.5pt;font-family:Consolas;color:black">->GetOutput()->GetSpacing()[i];</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">if</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> ((inputSize[i] != outputSize[i]) || (inputSpacing[i] != outputSpacing[i]))</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> {</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> std::cerr
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan"><<</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:#A31515">"The CUDA laplacian filter can only handle input and output regions of equal size and spacing"</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan"><<</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> std::</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">endl</span><span style="font-size:9.5pt;font-family:Consolas;color:black">;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">exit</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(1);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> }</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> }</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *pin = *(</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black">**)(</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">this</span><span style="font-size:9.5pt;font-family:Consolas;color:black">->GetInput()->GetCudaDataManager()->GetGPUBufferPointer());</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> // This is a test area</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto;margin-left:36.0pt;text-indent:36.0pt">
<span style="font-size:9.5pt;font-family:Consolas;color:blue">typename</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">InputImageType</span><span style="font-size:9.5pt;font-family:Consolas;color:black">::</span><span style="font-size:9.5pt;font-family:Consolas;color:darkolivegreen">IndexType</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
index;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> index.Fill(0);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">typename</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">InputImageType</span><span style="font-size:9.5pt;font-family:Consolas;color:black">::</span><span style="font-size:9.5pt;font-family:Consolas;color:darkolivegreen">SizeType</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
size;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">for</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> (</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">auto</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
i = 0; i < 3; ++i)</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> size.Fill(</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">this</span><span style="font-size:9.5pt;font-family:Consolas;color:black">->GetInput()->GetLargestPossibleRegion().GetSize()[i]);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">typename</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">InputImageType</span><span style="font-size:9.5pt;font-family:Consolas;color:black">::</span><span style="font-size:9.5pt;font-family:Consolas;color:darkolivegreen">RegionType</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
region(index, size);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> // images for gradients</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:blue">auto</span><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:black"> grad_x
</span><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:darkcyan">=</span><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:darkcyan">CudaImage</span><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:black"><</span><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:black">,
3>::New();</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:black"> grad_x</span><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:black">SetRegions(region);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:black">grad_x</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">Allocate();</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> grad_x</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">FillBuffer(1);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">auto</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> grad_y
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">=</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">CudaImage</span><span style="font-size:9.5pt;font-family:Consolas;color:black"><</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
3>::New();</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> grad_y</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">SetRegions(region);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> grad_y</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">Allocate();</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">auto</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> grad_z
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">=</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">CudaImage</span><span style="font-size:9.5pt;font-family:Consolas;color:black"><</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
3>::New();</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> grad_z</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">SetRegions(region);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> grad_z</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">Allocate();</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *pout_x = *(</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black">**)(grad_x</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">GetCudaDataManager()</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">GetGPUBufferPointer());</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *pout_y = *(</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black">**)(grad_y</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">GetCudaDataManager()</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">GetGPUBufferPointer());</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *pout_z = *(</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black">**)(grad_z</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">GetCudaDataManager()</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">GetGPUBufferPointer());</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">CUDA_gradient</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(inputSize, inputSpacing, pin, pout_x, pout_y, pout_z); // after this line neither of the pout_(xyz)
images have changed.</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> // put the gradient images in a single covariant vector image</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">auto</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> CompositeImageFilter
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">=</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> itk::</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">ComposeImageFilter</span><span style="font-size:9.5pt;font-family:Consolas;color:black"><</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">CudaImage</span><span style="font-size:9.5pt;font-family:Consolas;color:black"><</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
3>, </span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">CudaImage</span><span style="font-size:9.5pt;font-family:Consolas;color:black"><</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">CovariantVector</span><span style="font-size:9.5pt;font-family:Consolas;color:black"><</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
3>,3>>::New();</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> CompositeImageFilter</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">SetInput1(grad_x);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> CompositeImageFilter</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">SetInput2(grad_y);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> CompositeImageFilter</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">SetInput3(grad_z);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> CompositeImageFilter</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">Update();</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">this</span><span style="font-size:9.5pt;font-family:Consolas;color:black">->GetOutput()->Graft(CompositeImageFilter</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">GetOutput());</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> }</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="color:black">The cuda/kernel function</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:#6F008A">__global__</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">void</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">gradient_kernel</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
* </span><span style="font-size:9.5pt;font-family:Consolas;color:gray">in</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">grad_x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">grad_y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">grad_z</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">int3</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">float3</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Spacing</span><span style="font-size:9.5pt;font-family:Consolas;color:black">);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:blue">void</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">CUDA_gradient</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[3],</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">spacing</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[3],</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dev_in</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dev_out_x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dev_out_y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dev_out_z</span><span style="font-size:9.5pt;font-family:Consolas;color:black">)</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">{</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">int3</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> dev_Size
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">=</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">make_int3</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[0],
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[1],
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[2]);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">float3</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> dev_Spacing
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">=</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">make_float3</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">spacing</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[0],
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">spacing</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[1],
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">spacing</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[2]);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:green">// Output volume</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">long</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> outputMemorySize =
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[0] *
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[1] *
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[2] *
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">sizeof</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black">);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaMalloc</span><span style="font-size:9.5pt;font-family:Consolas;color:black">((</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">void</span><span style="font-size:9.5pt;font-family:Consolas;color:black">**)&</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dev_out_x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
outputMemorySize);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaMalloc</span><span style="font-size:9.5pt;font-family:Consolas;color:black">((</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">void</span><span style="font-size:9.5pt;font-family:Consolas;color:black">**)&</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dev_out_y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
outputMemorySize);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaMalloc</span><span style="font-size:9.5pt;font-family:Consolas;color:black">((</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">void</span><span style="font-size:9.5pt;font-family:Consolas;color:black">**)&</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dev_out_z</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
outputMemorySize);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaMemset</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dev_out_x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
0, outputMemorySize);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaMemset</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dev_out_y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
0, outputMemorySize);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaMemset</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dev_out_z</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
0, outputMemorySize);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">printf</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:#A31515">"Device Variable Copying:\t</span><span style="font-size:9.5pt;font-family:Consolas;color:mediumseagreen">%s</span><span style="font-size:9.5pt;font-family:Consolas;color:#A31515">\n"</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaGetErrorString</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaGetLastError</span><span style="font-size:9.5pt;font-family:Consolas;color:black">()));</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:green">// Thread Block Dimensions</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">dim3</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> dimBlock
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">=</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">dim3(</span><span style="font-size:9.5pt;font-family:Consolas;color:black">16, 4, 4</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">)</span><span style="font-size:9.5pt;font-family:Consolas;color:black">;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> blocksInX =
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">iDivUp</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[0],
dimBlock.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> blocksInY =
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">iDivUp</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[1],
dimBlock.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> blocksInZ =
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">iDivUp</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[2],
dimBlock.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">z</span><span style="font-size:9.5pt;font-family:Consolas;color:black">);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">dim3</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> dimGrid
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">=</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">dim3(</span><span style="font-size:9.5pt;font-family:Consolas;color:black">blocksInX, blocksInY, blocksInZ</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">)</span><span style="font-size:9.5pt;font-family:Consolas;color:black">;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">gradient_kernel</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> <<< dimGrid, dimBlock >>> (dev_in, dev_out_x, dev_out_y, dev_out_z, dev_Size, dev_Spacing);</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaDeviceSynchronize</span><span style="font-size:9.5pt;font-family:Consolas;color:black">();</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">printf</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:#A31515">"Device Variable Copying:\t</span><span style="font-size:9.5pt;font-family:Consolas;color:mediumseagreen">%s</span><span style="font-size:9.5pt;font-family:Consolas;color:#A31515">\n"</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaGetErrorString</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaGetLastError</span><span style="font-size:9.5pt;font-family:Consolas;color:black">()));</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:#6F008A">CUDA_CHECK_ERROR</span><span style="font-size:9.5pt;font-family:Consolas;color:black">;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">}</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:#6F008A">__global__</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:blue">void</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">gradient_kernel</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
* </span><span style="font-size:9.5pt;font-family:Consolas;color:gray">in</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">grad_x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">grad_y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">grad_z</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">int3</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">float3</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Spacing</span><span style="font-size:9.5pt;font-family:Consolas;color:black">)</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">{</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">unsigned</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> i = blockIdx.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
* blockDim.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">x</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> + threadIdx.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">unsigned</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> j = blockIdx.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
* blockDim.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">y</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> + threadIdx.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">unsigned</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> k = blockIdx.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">z</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
* blockDim.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">z</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> + threadIdx.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">z</span><span style="font-size:9.5pt;font-family:Consolas;color:black">;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">if</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> (i >=
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
|| j >= </span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
|| k >= </span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">z</span><span style="font-size:9.5pt;font-family:Consolas;color:black">)</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">return</span><span style="font-size:9.5pt;font-family:Consolas;color:black">;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">long</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> id = (k *
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
+ j) * </span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
+ i;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">long</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> id_x = (k *
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
+ j) * </span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
+ i + 1;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">long</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> id_y = (k *
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
+ j + 1)* </span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
+ i;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">long</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> id_z = ((k + 1) *
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
+ j) * </span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
+ i;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">if</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> (i == (</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
- 1)) </span><span style="font-size:9.5pt;font-family:Consolas;color:gray">grad_x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[id] = 0;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">else</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">grad_x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[id] = (</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">in</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[id_x]
- </span><span style="font-size:9.5pt;font-family:Consolas;color:gray">in</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[id]) /
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Spacing</span><span style="font-size:9.5pt;font-family:Consolas;color:black">.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">x</span><span style="font-size:9.5pt;font-family:Consolas;color:black">;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">if</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> (j == (</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
- 1)) </span><span style="font-size:9.5pt;font-family:Consolas;color:gray">grad_y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[id] = 0;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">else</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">grad_y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[id] = (</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">in</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[id_y]
- </span><span style="font-size:9.5pt;font-family:Consolas;color:gray">in</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[id]) /
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Spacing</span><span style="font-size:9.5pt;font-family:Consolas;color:black">.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">y</span><span style="font-size:9.5pt;font-family:Consolas;color:black">;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">if</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> (k == (</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">z</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
- 1)) </span><span style="font-size:9.5pt;font-family:Consolas;color:gray">grad_z</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[id] = 0;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">else</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">grad_z</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[id] = (</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">in</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[id_z]
- </span><span style="font-size:9.5pt;font-family:Consolas;color:gray">in</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[id]) /
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">c_Spacing</span><span style="font-size:9.5pt;font-family:Consolas;color:black">.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">z</span><span style="font-size:9.5pt;font-family:Consolas;color:black">;</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="font-size:9.5pt;font-family:Consolas;color:black">}</span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"><span style="color:black"> </span><o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:auto;mso-margin-bottom-alt:auto"> <o:p></o:p></p>
</div>
</div>
<p class="MsoNormal">_______________________________________________<br>
Rtk-users mailing list<br>
<a href="mailto:Rtk-users@public.kitware.com" target="_blank">Rtk-users@public.kitware.com</a><br>
<a href="https://public.kitware.com/mailman/listinfo/rtk-users" target="_blank">https://public.kitware.com/mailman/listinfo/rtk-users</a><o:p></o:p></p>
</blockquote>
</div>
</div>
</body>
</html>