<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;}
@font-face
{font-family:Verdana;
panose-1:2 11 6 4 3 5 4 4 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">Hello,<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 want to provide a solution to the question I had. In the case that someone else need an code example.<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1F497D">Just to give an idea how the implementation can look like.
<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1F497D">The main reads an image and passes this image to the function that calls the kernel.<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1F497D">After the kernel is done the result is copied to the new vector image [1].<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">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">main</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">
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">argc</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">char</span><span style="font-size:9.5pt;font-family:Consolas;color:black">*
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">argv</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:darkslateblue">createImage</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">auto</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> nrrdIO
</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">NRRDIOType</span><span style="font-size:9.5pt;font-family:Consolas;color:black">::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">
</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"> reader
</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">ReaderType</span><span style="font-size:9.5pt;font-family:Consolas;color:black">::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"> reader</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">SetFileName(</span><span style="font-size:9.5pt;font-family:Consolas;color:#A31515">"newImage.nrrd"</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"> reader</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">SetImageIO(nrrdIO);<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">try</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"> reader</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">catch</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">ExceptionObject</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
e)<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">"Exception caught!\n"</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:black">std::cerr
</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"> e
</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"> std::</span><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">endl</span><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:black">;<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">}<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">auto</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> image = reader</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" 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">auto</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> 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"> image->GetLargestPossibleRegion().GetSize();<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"> sz[3] = {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">0</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">]</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,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">1</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">]</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,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">2</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">
</span><span lang="PL" style="font-size:9.5pt;font-family:Consolas;color:blue">auto</span><span lang="PL" style="font-size:9.5pt;font-family:Consolas;color:black"> totalVoxel = sz[0] * sz[1] * sz[2];<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span lang="PL" style="font-size:9.5pt;font-family:Consolas;color:black"><o:p> </o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span lang="PL" 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"> cuda
</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">CudaTest()</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">auto</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> out =
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">new</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">float4</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[totalVoxel];<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> cuda.runKernel(out, image->GetBufferPointer(), sz);<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">auto</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> vecImage
</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">VectorImageType</span><span style="font-size:9.5pt;font-family:Consolas;color:black">::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">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">VectorImageType</span><span style="font-size:9.5pt;font-family:Consolas;color:black">::</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">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</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">[</span><span style="font-size:9.5pt;font-family:Consolas;color:black">0</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">]</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
= 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"> index</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">[</span><span style="font-size:9.5pt;font-family:Consolas;color:black">1</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">]</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
= 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"> index</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">[</span><span style="font-size:9.5pt;font-family:Consolas;color:black">2</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">]</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
= 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:darkcyan">VectorImageType</span><span style="font-size:9.5pt;font-family:Consolas;color:black">::</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">RegionType</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
region</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">(</span><span style="font-size:9.5pt;font-family:Consolas;color:black">index,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">;<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> vecImage</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"> vecImage</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">SetVectorLength(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"> vecImage</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:green">// copy image buffer to vecImage, component by component</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">auto</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> vecBuffer = vecImage</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">GetBufferPointer();<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"> j = 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">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 < totalVoxel; ++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"> vecBuffer[j] = out[i].</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++;<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> vecBuffer[j] = out[i].</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++;<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> vecBuffer[j] = out[i].</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">;
j++;<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">auto</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> writer
</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">WriterType</span><span style="font-size:9.5pt;font-family:Consolas;color:black">::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"> writer</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">SetInput(vecImage);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> writer</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">SetFileName(</span><span style="font-size:9.5pt;font-family:Consolas;color:#A31515">"out.nrrd"</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"> writer</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">-></span><span style="font-size:9.5pt;font-family:Consolas;color:black">SetImageIO(nrrdIO);<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">try</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"> writer</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">catch</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">ExceptionObject</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
e)<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">"Exception caught!\n"</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:black">std::cerr
</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"> e
</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"> std::</span><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">endl</span><span lang="DE" style="font-size:9.5pt;font-family:Consolas;color:black">;<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">}<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">delete[]</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> out;<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="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">The function that starts the kernel looks like this:<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" 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">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">CudaTest</span><span style="font-size:9.5pt;font-family:Consolas;color:black">::runKernel(</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">float4</span><span style="font-size:9.5pt;font-family:Consolas;color:black">*
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">out</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">data</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">*
</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">)<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">auto</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> totalVoxel =
</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:green">// copy image to texture memory</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">auto</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> channelDesc
</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">cudaCreateChannelDesc</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:darkcyan">cudaArray</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *cuArray;<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"> extent
</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_cudaExtent</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:#6F008A">gpuErrChk</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaMalloc3DArray</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(&cuArray,
&channelDesc, extent));<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">cudaMemcpy3DParms</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> copyparms = { 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"> copyparms.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">extent</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"> extent;<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> copyparms.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">dstArray</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
= cuArray;<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> copyparms.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">kind</span><span style="font-size:9.5pt;font-family:Consolas;color:black">
= </span><span style="font-size:9.5pt;font-family:Consolas;color:#314F4F">cudaMemcpyHostToDevice</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"> copyparms.</span><span style="font-size:9.5pt;font-family:Consolas;color:darkred">srcPtr</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">
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">make_cudaPitchedPtr</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">data</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: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><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]);<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">gpuErrChk</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaMemcpy3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(©parms));<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">// set texture configuration and bind array to texture</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"> tex_volume.addressMode[0] =
</span><span style="font-size:9.5pt;font-family:Consolas;color:#314F4F">cudaAddressModeClamp</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"> tex_volume.addressMode[1] =
</span><span style="font-size:9.5pt;font-family:Consolas;color:#314F4F">cudaAddressModeClamp</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"> tex_volume.addressMode[2] =
</span><span style="font-size:9.5pt;font-family:Consolas;color:#314F4F">cudaAddressModeClamp</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"> tex_volume.filterMode =
</span><span style="font-size:9.5pt;font-family:Consolas;color:#314F4F">cudaFilterModeLinear</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"> tex_volume.normalized =
</span><span style="font-size:9.5pt;font-family:Consolas;color:blue">false</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">gpuErrChk</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaBindTextureToArray</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume,
cuArray, channelDesc));<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">// allocate memory for gradient image data on the device</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">float4</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> *dev_gradientImage;<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">gpuErrChk</span><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">(&dev_gradientImage,
totalVoxel * </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:darkcyan">float4</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">// copy image dimension to device</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">* dev_dimension;<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">gpuErrChk</span><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">(&dev_dimension,
3 * </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">int</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">gpuErrChk</span><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">(dev_dimension,
</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><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">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black">),
</span><span style="font-size:9.5pt;font-family:Consolas;color:#314F4F">cudaMemcpyHostToDevice</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">// start computation for the gradient image</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"> blockDim</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">(</span><span style="font-size:9.5pt;font-family:Consolas;color:black">16,
8, 8</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:green">//<- threads in block, change according to used GPU, max 1024 threads in a single block</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"> grid</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:gray">size</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[0]
/ 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">) + 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">[1]
/ 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">) + 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]
/ 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">) + 1</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">computeGradientImage</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> << <grid, blockDim >> >(dev_gradientImage, dev_dimension);<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">gpuErrChk</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">gpuErrChk</span><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"><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">// unbind texture</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">gpuErrChk</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaUnbindTexture</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume));<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">gpuErrChk</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">cudaFreeArray</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(cuArray));<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:#6F008A">gpuErrChk</span><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">out</span><span style="font-size:9.5pt;font-family:Consolas;color:black">,
dev_gradientImage, totalVoxel * </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:darkcyan">float4</span><span style="font-size:9.5pt;font-family:Consolas;color:black">),
</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: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">"Kernel is done.\n"</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">The kernel used in this example.<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:#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">computeGradientImage</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(</span><span style="font-size:9.5pt;font-family:Consolas;color:darkcyan">float4</span><span style="font-size:9.5pt;font-family:Consolas;color:black">*
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">gradientImage</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">*
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dimension</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">// every thread computes the float4 voxel with theta,phi,magnitude from gradient image</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"> idx = 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">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> idy = 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">int</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> idz = 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"> (idx <
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dimension</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[0] && idy <
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dimension</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[1] && idz <
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dimension</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">// define sobel filter for each direction</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"> x_000 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx - 1, idy - 1, idz - 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">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> x_001 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx - 1, idy - 1, idz );<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"> x_002 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx - 1, idy - 1, idz + 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">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> x_010 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx - 1, idy, idz - 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">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> x_011 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx - 1, idy, idz );<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"> x_012 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx - 1, idy, idz + 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">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> x_020 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx - 1, idy + 1, idz - 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">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> x_021 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx - 1, idy + 1, idz );<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"> x_022 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx - 1, idy + 1, idz + 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">
</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"> x_100 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx, idy - 1, idz - 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">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> x_101 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx, idy - 1, idz );<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"> x_102 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx, idy - 1, idz + 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">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> x_110 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx, idy, idz - 1);<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:green">// float x_111 = tex3D(tex_volume, idx, idy, idz );</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"> x_112 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx, idy, idz + 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">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> x_120 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx, idy + 1, idz - 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">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> x_121 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx, idy + 1, idz );<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"> x_122 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx, idy + 1, idz + 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">
</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"> x_200 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx + 1, idy - 1, idz - 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">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> x_201 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx + 1, idy - 1, idz );<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"> x_202 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx + 1, idy - 1, idz + 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">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> x_210 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx + 1, idy, idz - 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">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> x_211 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx + 1, idy, idz );<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"> x_212 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx + 1, idy, idz + 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">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> x_220 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx + 1, idy + 1, idz - 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">float</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> x_221 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx + 1, idy + 1, idz );<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"> x_222 =
</span><span style="font-size:9.5pt;font-family:Consolas;color:red">tex3D</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume, idx + 1, idy + 1, idz + 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">
</span><span style="font-size:9.5pt;font-family:Consolas;color:green">// derivative in x direction</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">auto</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> centerX = - 1.f*x_000 - 3.f*x_010 - 1.f*x_020<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> - 3.f*x_001 - 6.f*x_011 - 3.f*x_021<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> - 1.f*x_002 - 3.f*x_012 - 1.f*x_022<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> + 1.f*x_200 + 3.f*x_210 + 1.f*x_220<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> + 3.f*x_201 + 6.f*x_211 + 3.f*x_221<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> + 1.f*x_202 + 3.f*x_212 + 1.f*x_222;<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">// derivative in y direction</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">auto</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> centerY = - 1.f*x_000 - 3.f*x_100 - 1.f*x_200<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> - 3.f*x_001 - 6.f*x_101 - 3.f*x_201<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> - 1.f*x_002 - 3.f*x_102 - 1.f*x_202<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> + 1.f*x_020 + 3.f*x_120 + 1.f*x_220<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> + 3.f*x_021 + 6.f*x_121 + 3.f*x_221<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> + 1.f*x_022 + 3.f*x_122 + 1.f*x_222;<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">// derivative in z direction</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">auto</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> centerZ = - 1.f*x_000 - 3.f*x_100 - 1.f*x_200<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> - 3.f*x_010 - 6.f*x_110 - 3.f*x_210<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> - 1.f*x_020 - 3.f*x_120 - 1.f*x_220<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> + 1.f*x_002 + 3.f*x_102 + 1.f*x_202<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> + 3.f*x_012 + 6.f*x_112 + 3.f*x_212<o:p></o:p></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> + 1.f*x_022 + 3.f*x_122 + 1.f*x_222;<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">// magnitude of each voxels 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">auto</span><span style="font-size:9.5pt;font-family:Consolas;color:black"> magn =
</span><span style="font-size:9.5pt;font-family:Consolas;color:darkslateblue">sqrtf</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(centerX*centerX + centerY*centerY + centerZ*centerZ);<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:gray">gradientImage</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[idx +
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dimension</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[0] * (idy + idz *
</span><span style="font-size:9.5pt;font-family:Consolas;color:gray">dimension</span><span style="font-size:9.5pt;font-family:Consolas;color:black">[1])]
</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_float4</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(centerX, centerY, centerZ, magn);<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="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">with kind regards,<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:9.5pt;font-family:Consolas;color:black">Gordian</span><o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">[1]:<span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1F497D">
<a href="https://itk.org/Doxygen/html/classitk_1_1VectorImage.html">https://itk.org/Doxygen/html/classitk_1_1VectorImage.html</a><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"><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""> Dženan Zukić [mailto:dzenanz@gmail.com]
<br>
<b>Gesendet:</b> Mittwoch, 28. Juni 2017 17:17<br>
<b>An:</b> Kabelitz, Gordian<br>
<b>Cc:</b> insight-users@itk.org<br>
<b>Betreff:</b> Re: [ITK-users] Writing from an external buffer to VectorImage<o:p></o:p></span></p>
<p class="MsoNormal"><o:p> </o:p></p>
<div>
<div>
<p class="MsoNormal"><span style="font-family:"Verdana","sans-serif"">Hi Gordian,<o:p></o:p></span></p>
</div>
<div>
<p class="MsoNormal"><span style="font-family:"Verdana","sans-serif""><o:p> </o:p></span></p>
</div>
<div>
<p class="MsoNormal"><span style="font-family:"Verdana","sans-serif"">this approach looks like it should work. What is wrong with it?<o:p></o:p></span></p>
</div>
<div>
<div>
<div>
<p class="MsoNormal"><span style="font-family:"Verdana","sans-serif""><o:p> </o:p></span></p>
</div>
<div>
<p class="MsoNormal"><span style="font-family:"Verdana","sans-serif"">Regards,<o:p></o:p></span></p>
</div>
<div>
<p class="MsoNormal"><span style="font-family:"Verdana","sans-serif"">Dženan Zukić, PhD, Senior R&D Engineer, Kitware (Carrboro, N.C.)<o:p></o:p></span></p>
</div>
</div>
</div>
</div>
<div>
<p class="MsoNormal"><o:p> </o:p></p>
<div>
<p class="MsoNormal">On Wed, Jun 28, 2017 at 9:51 AM, Kabelitz, Gordian <<a href="mailto:Gordian.Kabelitz@medma.uni-heidelberg.de" target="_blank">Gordian.Kabelitz@medma.uni-heidelberg.de</a>> wrote:<o:p></o:p></p>
<p class="MsoNormal">Hello,<br>
<br>
i computed a gradient with my own function and as a result a pointer to an image buffer is provided. I know the size, origin and spacing of the gradient component image.<br>
I want to copy the gradient image into an itk::VectorImage with the components for the x,y,z gradients.<br>
<br>
The way I copied the image to the GPU is that I retrieved the buffer pointer from my input image and use the pointer to copy the image data to the GPU.<br>
I used the way proposed in [1]. The computeGradientImage method is listed at the end of this mail.<br>
<br>
[...]<br>
// get float pointer to image data<br>
ImageType::Pointer image = reader->GetOutput();<br>
Image->Update();<br>
<br>
float* data = image.GetBufferPointer();<br>
// copy image data to GPU texture memory (this works)<br>
gpu_dev->setVoxels(dimension, voxelSize, data);<br>
[...]<br>
computeGradientImage<<<n,m>>> (dev_gradientImage, dimension);<br>
<br>
// copy resulting gradientImage to host variable<br>
float4* host_gradientImage;<br>
cudaMemcpy(host_gradient, dev_gradientImage, numberOfVoxels*sizeof(float4));<br>
<br>
--> Pseudo Code <--<br>
// Now I want to reverse the copy process. I have a float4 image and want to copy this into a itk::VectorImage with VariableVectorLength of 3 (skipping the magnitude value).<br>
[...] -> size, spacing, origin, region definition<br>
Itk::VectorImageType vecImage = VectorImageType::New();<br>
vecImage->setRegion(region);<br>
vecImage ->SetVectorLength(3);<br>
vecImage->Allocate();<br>
<br>
// copy image buffer to vecImage, component by component<br>
auto vecBuffer = vecImage->getBufferPointer();<br>
auto j = 0;<br>
for (i=0; i<totalVoxel; ++i)<br>
{<br>
vecbuffer[j] = host_gradient[i].x; j++;<br>
vecbuffer[j] = host_gradient[i].y; j++;<br>
vecbuffer[j] = host_gradient[i].z; j++;<br>
}<br>
<br>
// save vecImage as nrrd image<br>
[...]<br>
<br>
I haven't found a way to achieve my idea.<br>
Are there any suggestions or examples?<br>
As far I can see I cannot use the itk::ImportImageFilter.<br>
<br>
Thank you for any suggestions.<br>
With kind regards,<br>
Gordian<br>
<br>
[1]: <a href="https://itk.org/CourseWare/Training/GettingStarted-V.pdf" target="_blank">
https://itk.org/CourseWare/Training/GettingStarted-V.pdf</a><br>
<br>
void computeGradientImage(float4* gradientImage, int* dimension)<br>
{<br>
// every thread computes the float4 voxel with theta,phi,magnitude from gradient image<br>
int idx = blockIdx.x * blockDim.x + threadIdx.x;<br>
int idy = blockIdx.y * blockDim.y + threadIdx.y;<br>
int idz = blockIdx.z * blockDim.z + threadIdx.z;<br>
<br>
if (idx < dimension[0] && idy < dimension[1] && idz < dimension[2])<br>
{<br>
// define sobel filter for each direction<br>
[...]<br>
<br>
// run sobel on image in texture memory for each direction and put result into a float4 image<br>
gradientImage[idx + dimension[0] * (idy + idz * dimension[1])] = make_float4(sobelX, sobelY, sobelZ, magn);<br>
}<br>
}<br>
<br>
<br>
_____________________________________<br>
Powered by <a href="http://www.kitware.com" target="_blank">www.kitware.com</a><br>
<br>
Visit other Kitware open-source projects at<br>
<a href="http://www.kitware.com/opensource/opensource.html" target="_blank">http://www.kitware.com/opensource/opensource.html</a><br>
<br>
Kitware offers ITK Training Courses, for more information visit:<br>
<a href="http://www.kitware.com/products/protraining.php" target="_blank">http://www.kitware.com/products/protraining.php</a><br>
<br>
Please keep messages on-topic and check the ITK FAQ at:<br>
<a href="http://www.itk.org/Wiki/ITK_FAQ" target="_blank">http://www.itk.org/Wiki/ITK_FAQ</a><br>
<br>
Follow this link to subscribe/unsubscribe:<br>
<a href="http://public.kitware.com/mailman/listinfo/insight-users" target="_blank">http://public.kitware.com/mailman/listinfo/insight-users</a><o:p></o:p></p>
</div>
<p class="MsoNormal"><o:p> </o:p></p>
</div>
</div>
</body>
</html>