<div dir="ltr"><div class="gmail_default" style="font-family:verdana,sans-serif;font-size:small">Thanks for sharing your code Gordian, it will probably help somebody in the future!</div></div><div class="gmail_extra"><br><div class="gmail_quote">On Mon, Jul 3, 2017 at 7:22 AM, Kabelitz, Gordian <span dir="ltr"><<a href="mailto:Gordian.Kabelitz@medma.uni-heidelberg.de" target="_blank">Gordian.Kabelitz@medma.uni-heidelberg.de</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<div lang="EN-US" link="blue" vlink="purple">
<div class="m_5190462275529389030WordSection1">
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d">Hello,<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d"><u></u> <u></u></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.<u></u><u></u></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.
<u></u><u></u></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.<u></u><u></u></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].<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d"><u></u> <u></u></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">[])<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">{<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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">();<u></u><u></u></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();<u></u><u></u></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();<u></u><u></u></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.<wbr>nrrd"</span><span style="font-size:9.5pt;font-family:Consolas;color:black">);<u></u><u></u></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);<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
<u></u><u></u></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"><u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> {<u></u><u></u></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();<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> }<u></u><u></u></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)<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> {<u></u><u></u></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">;<u></u><u></u></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">;<u></u><u></u></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">}<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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();<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
<u></u><u></u></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-><wbr>GetLargestPossibleRegion().<wbr>GetSize();<u></u><u></u></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">};<u></u><u></u></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];<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span lang="PL" style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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">;<u></u><u></u></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];<u></u><u></u></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);<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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();<u></u><u></u></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;<u></u><u></u></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;<u></u><u></u></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;<u></u><u></u></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;<u></u><u></u></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">;<u></u><u></u></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);<u></u><u></u></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);<u></u><u></u></span></p><span class="">
<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();<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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"><u></u><u></u></span></p>
</span><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();<u></u><u></u></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;<u></u><u></u></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)<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> {<u></u><u></u></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++;<u></u><u></u></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++;<u></u><u></u></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++;<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> }<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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();<u></u><u></u></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);<u></u><u></u></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"><wbr>);<u></u><u></u></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);<u></u><u></u></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"><u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> {<u></u><u></u></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();<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> }<u></u><u></u></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)<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> {<u></u><u></u></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">;<u></u><u></u></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">;<u></u><u></u></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">}<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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;<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">
<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:9.5pt;font-family:Consolas;color:black">}<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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:<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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">)<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">{<u></u><u></u></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];<u></u><u></u></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"><u></u><u></u></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">>()<wbr>;<u></u><u></u></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;<u></u><u></u></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]);<u></u><u></u></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">(&<wbr>cuArray,
&channelDesc, extent));<u></u><u></u></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 };<u></u><u></u></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;<u></u><u></u></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;<u></u><u></u></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">;<u></u><u></u></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">dat<wbr>a</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]);<u></u><u></u></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">(&<wbr>copyparms));<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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"><u></u><u></u></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">;<u></u><u></u></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">;<u></u><u></u></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">;<u></u><u></u></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">;<u></u><u></u></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">;<u></u><u></u></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">cudaBindTextureToArr<wbr>ay</span><span style="font-size:9.5pt;font-family:Consolas;color:black">(tex_volume,
cuArray, channelDesc));<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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"><u></u><u></u></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;<u></u><u></u></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_<wbr>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">)));<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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"><u></u><u></u></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;<u></u><u></u></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_<wbr>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">)));<u></u><u></u></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_<wbr>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">));<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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"><u></u><u></u></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"><u></u><u></u></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">;<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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);<u></u><u></u></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">());<u></u><u></u></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">cudaDeviceSynchroniz<wbr>e</span><span style="font-size:9.5pt;font-family:Consolas;color:black">());<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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"><u></u><u></u></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">(<wbr>tex_volume));<u></u><u></u></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">(<wbr>cuArray));<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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">));<u></u><u></u></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">);<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:9.5pt;font-family:Consolas;color:black">}<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:black">The kernel used in this example.<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d"><u></u> <u></u></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">)<u></u><u></u></span></p><span class="">
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black">{<u></u><u></u></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"><u></u><u></u></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">;<u></u><u></u></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">;<u></u><u></u></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">;<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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])<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> {<u></u><u></u></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"><u></u><u></u></span></p>
</span><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);<u></u><u></u></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 );<u></u><u></u></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);<u></u><u></u></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);<u></u><u></u></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 );<u></u><u></u></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);<u></u><u></u></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);<u></u><u></u></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 );<u></u><u></u></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);<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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);<u></u><u></u></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 );<u></u><u></u></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);<u></u><u></u></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);<u></u><u></u></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"><u></u><u></u></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);<u></u><u></u></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);<u></u><u></u></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 );<u></u><u></u></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);<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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);<u></u><u></u></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 );<u></u><u></u></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);<u></u><u></u></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);<u></u><u></u></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 );<u></u><u></u></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);<u></u><u></u></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);<u></u><u></u></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 );<u></u><u></u></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);<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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"><u></u><u></u></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<u></u><u></u></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<u></u><u></u></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<u></u><u></u></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<u></u><u></u></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<u></u><u></u></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;<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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"><u></u><u></u></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<u></u><u></u></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<u></u><u></u></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<u></u><u></u></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<u></u><u></u></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<u></u><u></u></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;<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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"><u></u><u></u></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<u></u><u></u></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<u></u><u></u></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<u></u><u></u></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<u></u><u></u></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<u></u><u></u></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;<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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"><u></u><u></u></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);<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></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);<u></u><u></u></span></p>
<p class="MsoNormal" style="text-autospace:none"><span style="font-size:9.5pt;font-family:Consolas;color:black"> }<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:9.5pt;font-family:Consolas;color:black">}<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:9.5pt;font-family:Consolas;color:black"><u></u> <u></u></span></p>
<p class="MsoNormal"><span style="font-size:9.5pt;font-family:Consolas;color:black">with kind regards,<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:9.5pt;font-family:Consolas;color:black">Gordian</span><u></u><u></u></p>
<p class="MsoNormal"><u></u> <u></u></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" target="_blank">https://itk.org/Doxygen/html/<wbr>classitk_1_1VectorImage.html</a><u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d"><u></u> <u></u></span></p>
<p class="MsoNormal"><u></u> <u></u></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:<a href="mailto:dzenanz@gmail.com" target="_blank">dzenanz@gmail.com</a>]
<br>
<b>Gesendet:</b> Mittwoch, 28. Juni 2017 17:17<br>
<b>An:</b> Kabelitz, Gordian<br>
<b>Cc:</b> <a href="mailto:insight-users@itk.org" target="_blank">insight-users@itk.org</a><br>
<b>Betreff:</b> Re: [ITK-users] Writing from an external buffer to VectorImage<u></u><u></u></span></p><div><div class="h5">
<p class="MsoNormal"><u></u> <u></u></p>
<div>
<div>
<p class="MsoNormal"><span style="font-family:"Verdana","sans-serif"">Hi Gordian,<u></u><u></u></span></p>
</div>
<div>
<p class="MsoNormal"><span style="font-family:"Verdana","sans-serif""><u></u> <u></u></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?<u></u><u></u></span></p>
</div>
<div>
<div>
<div>
<p class="MsoNormal"><span style="font-family:"Verdana","sans-serif""><u></u> <u></u></span></p>
</div>
<div>
<p class="MsoNormal"><span style="font-family:"Verdana","sans-serif"">Regards,<u></u><u></u></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.)<u></u><u></u></span></p>
</div>
</div>
</div>
</div>
<div>
<p class="MsoNormal"><u></u> <u></u></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-<wbr>heidelberg.de</a>> wrote:<u></u><u></u></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))<wbr>;<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/<wbr>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>
______________________________<wbr>_______<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/<wbr>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/<wbr>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_<wbr>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/<wbr>mailman/listinfo/insight-users</a><u></u><u></u></p>
</div>
<p class="MsoNormal"><u></u> <u></u></p>
</div>
</div></div></div>
</div>
<br>______________________________<wbr>_______<br>
Powered by <a href="http://www.kitware.com" rel="noreferrer" 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" rel="noreferrer" target="_blank">http://www.kitware.com/<wbr>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" rel="noreferrer" target="_blank">http://www.kitware.com/<wbr>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" rel="noreferrer" target="_blank">http://www.itk.org/Wiki/ITK_<wbr>FAQ</a><br>
<br>
Follow this link to subscribe/unsubscribe:<br>
<a href="http://public.kitware.com/mailman/listinfo/insight-users" rel="noreferrer" target="_blank">http://public.kitware.com/<wbr>mailman/listinfo/insight-users</a><br>
<br></blockquote></div><br></div>