<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 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: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:11.0pt;
font-family:"Calibri","sans-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-compose;
font-family:"Calibri","sans-serif";
color:windowtext;}
.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" style="mso-margin-top-alt:7.5pt;margin-right:0cm;margin-bottom:3.75pt;margin-left:0cm;line-height:12.0pt">
Hi rtk-users,<o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:7.5pt;margin-right:0cm;margin-bottom:3.75pt;margin-left:0cm;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:7.5pt;margin-right:0cm;margin-bottom:3.75pt;margin-left:0cm;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:7.5pt;margin-right:0cm;margin-bottom:3.75pt;margin-left:0cm;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:7.5pt;margin-right:0cm;margin-bottom:3.75pt;margin-left:0cm;line-height:12.0pt">
With kind regards,<o:p></o:p></p>
<p class="MsoNormal" style="mso-margin-top-alt:7.5pt;margin-right:0cm;margin-bottom:3.75pt;margin-left:0cm;line-height:12.0pt">
Gordian<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">The GPUGenerateData function:<o:p></o:p></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">GPUGenerateData()<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"> inputSize[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">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> outputSize[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"> inputSpacing[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"> outputSpacing[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"><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">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++)<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"> 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];<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><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];<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><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];<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><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];<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">if</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> ((inputSize[i] != outputSize[i]) || (inputSpacing[i] != outputSpacing[i]))<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"> 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">;<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">exit</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(1);<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"><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"> *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());<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"> // This is a test area<o:p></o:p></span></p>
<p class="MsoNormal" style="margin-left:36.0pt;text-indent:36.0pt;text-autospace:none">
<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;<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> index.Fill(0);<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">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;<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">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)<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><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]);<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">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);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> // images for gradients<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">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();<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"> 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);<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: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();<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><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);<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">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();<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><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);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><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();<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">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();<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><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);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><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();<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">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());<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"> *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());<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"> *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());<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">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.<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"> // put the gradient images in a single covariant vector image<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">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();<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><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);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><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);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><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);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><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();<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">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());<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="color:black">The cuda/kernel function<o:p></o:p></span></p>
<p class="MsoNormal"><span style="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:#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">);<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: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">
</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 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);<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">**)&</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);<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">**)&</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);<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">(</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);<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">(</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);<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">(</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);<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">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">()));<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">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">;<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_out_x, dev_out_y, dev_out_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:darkslateblue">cudaDeviceSynchronize</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">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">()));<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"><o:p> </o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:#6F008A">__global__</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: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">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">)<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: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">;<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">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">;<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">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">;<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">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">)<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">return</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">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;<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"> 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;<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"> 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;<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"> 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;<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">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;<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">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">;<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">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;<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">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">;<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">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;<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">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">;<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"><span style="color:black"><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="color:black"><o:p> </o:p></span></p>
<p class="MsoNormal"><o:p> </o:p></p>
</div>
</body>
</html>