<html><head></head><body style="word-wrap: break-word; -webkit-nbsp-mode: space; -webkit-line-break: after-white-space; ">Hi<div>As a hint i found these 2 writings:</div><div><br></div><div><span class="Apple-style-span" style="line-height: 24px; "><span style="margin-top: 0px; margin-right: 0px; margin-bottom: 0px; margin-left: 0px; padding-top: 0px; padding-right: 0px; padding-bottom: 0px; padding-left: 0px; border-top-width: 0px; border-right-width: 0px; border-bottom-width: 0px; border-left-width: 0px; border-style: initial; border-color: initial; background-image: initial; background-attachment: initial; background-origin: initial; background-clip: initial; background-color: transparent;">So about that <span class="simple_code" style="margin-top: 0px; margin-right: 0px; margin-bottom: 0px; margin-left: 0px; padding-top: 0px; padding-right: 0px; padding-bottom: 0px; padding-left: 0px; border-top-width: 0px; border-right-width: 0px; border-bottom-width: 0px; border-left-width: 0px; border-style: initial; border-color: initial; background-image: initial; background-attachment: initial; background-origin: initial; background-clip: initial; ">float padding[2];</span> in the struct definition. This is because of memory<span style="margin-top: 0px; margin-right: 0px; margin-bottom: 0px; margin-left: 0px; padding-top: 0px; padding-right: 0px; padding-bottom: 0px; padding-left: 0px; border-top-width: 0px; border-right-width: 0px; border-bottom-width: 0px; border-left-width: 0px; border-style: initial; border-color: initial; background-image: initial; background-attachment: initial; background-origin: initial; background-clip: initial; ">alignment</span> in OpenCL. The best explanation I’ve seen so far is by <a href="http://forums.amd.com/forum/messageview.cfm?catid=390&threadid=122209&forumid=9" style="margin-top: 0px; margin-right: 0px; margin-bottom: 0px; margin-left: 0px; padding-top: 0px; padding-right: 0px; padding-bottom: 0px; padding-left: 0px; border-top-width: 0px; border-right-width: 0px; border-bottom-width: 0px; border-left-width: 0px; border-style: initial; border-color: initial; background-image: initial; background-attachment: initial; background-origin: initial; background-clip: initial; ">AndreasStahl</a> which I will briefly summarize in relation to the struct above.<br>When interpreting a struct, OpenCL accesses the memory in blocks of 16 bytes, which is the same as 4 floats (each 4 bytes). So in our example if we did not have the padding, we would not be able to access our int because opencl would have interpreted it as the 3rd float out of the first 16bytes. This can get even more complicated if you have an array of structs, because then the size of you’re struct will need to be a multiple of 16, as explained in the linked forum post.</span></span></div><div><span class="Apple-style-span" style="line-height: 24px; "><span style="margin-top: 0px; margin-right: 0px; margin-bottom: 0px; margin-left: 0px; padding-top: 0px; padding-right: 0px; padding-bottom: 0px; padding-left: 0px; border-top-width: 0px; border-right-width: 0px; border-bottom-width: 0px; border-left-width: 0px; border-style: initial; border-color: initial; background-image: initial; background-attachment: initial; background-origin: initial; background-clip: initial; background-color: transparent;">The orig AMD posting:</span></span></div><div><span class="Apple-style-span" style="line-height: 24px; "><span style="margin-top: 0px; margin-right: 0px; margin-bottom: 0px; margin-left: 0px; padding-top: 0px; padding-right: 0px; padding-bottom: 0px; padding-left: 0px; border-top-width: 0px; border-right-width: 0px; border-bottom-width: 0px; border-left-width: 0px; border-style: initial; border-color: initial; background-image: initial; background-attachment: initial; background-origin: initial; background-clip: initial; background-color: transparent;"><br></span></span></div><div><span class="Apple-style-span" style="line-height: 24px; "><span style="margin-top: 0px; margin-right: 0px; margin-bottom: 0px; margin-left: 0px; padding-top: 0px; padding-right: 0px; padding-bottom: 0px; padding-left: 0px; border-top-width: 0px; border-right-width: 0px; border-bottom-width: 0px; border-left-width: 0px; border-style: initial; border-color: initial; background-image: initial; background-attachment: initial; background-origin: initial; background-clip: initial; background-color: transparent;"><a href="http://forums.amd.com/forum/messageview.cfm?catid=390&threadid=122209&forumid=9">http://forums.amd.com/forum/messageview.cfm?catid=390&threadid=122209&forumid=9</a></span></span></div><div><font class="Apple-style-span" color="#333333" face="Georgia, 'Bitstream Charter', serif"><span class="Apple-style-span" style="font-size: 16px; line-height: 24px; "><i><br></i></span></font></div><div>Conclusion: There are indeed places where the compiler does not take care of alignment ( bug? ), so we must do ourselfes.</div><div><br></div><div>Jens</div><div><br></div><div><br></div><div><br></div><div><br></div><div><br></div><div><br><div><div>Am 07.12.2011 um 16:38 schrieb Brecht Van Lommel:</div><br class="Apple-interchange-newline"><blockquote type="cite"><div>Hi,<br><br>Maybe it would be possible to set up an SSH login on a computer so I<br>can try and narrow this down further? There has to be a pattern here,<br>but I'm not sure what to suggest you to test, what I would do is to<br>keep removing code until it works, and find out exactly which kind of<br>struct variable layout or operation is causing issues.<br><br>On Wed, Dec 7, 2011 at 7:11 AM, Dalai Felinto <<a href="mailto:dfelinto@gmail.com">dfelinto@gmail.com</a>> wrote:<br><blockquote type="cite">"Good news", I believe both problems may indeed be related and of alignment.<br></blockquote><blockquote type="cite">At least the first one.<br></blockquote><blockquote type="cite"><br></blockquote><blockquote type="cite">If I change the alignment in kernel_type.h in simple ways [*] I get the Z<br></blockquote><blockquote type="cite">problem fixed, but the viewport broken in situations it was working before.<br></blockquote><blockquote type="cite">If I go to camera fly mode the viewport is fine ('dirty' but fine) until I<br></blockquote><blockquote type="cite">stop then I get the (2) problem. So what are the rules for alignment here?<br></blockquote><blockquote type="cite"><br></blockquote><blockquote type="cite">[*] <a href="http://www.pasteall.org/27124">http://www.pasteall.org/27124</a><br></blockquote><br>The alignment issue I was thinking of is when passing structs between<br>the C++ and OpenCL. If they layout variables in structs differently,<br>copying KernelData to the GPU would go wrong. So what I've tried to do<br>is align everything like we do in makesdna, and make all struct sizes<br>a multiple of 16 bytes.<br><br>What is quite strange, is that the struct where you added alignment<br>does not get passed between C++ and OpenCL, it stays purely in OpenCL,<br>and so I wouldn't expect this to be a problem. Maybe there's a<br>compiler bug related to nested structs?<br><br><blockquote type="cite">From the specification:<br></blockquote><br>"The OpenCL compiler is responsible for aligning data items to the<br>appropriate alignment as required by the data type. For arguments to a<br>__kernel function declared to be a pointer to a data type, the OpenCL<br>compiler can assume that the pointee is always appropriately aligned<br>as required by the data type."<br><br>So this means that in principle we shouldn't have to worry about<br>alignment of structs that stay on the GPU, and only for KernelData we<br>need to worry about alignment. But of course compiler bugs can make<br>this an issue.<br><br>"For 3-component vector data types, the size of the data type is 4 *<br>sizeof(component). This means that a 3-component vector data type will<br>be aligned to a 4 * sizeof(component) boundary."<br><br>This seems to indicate that differential3 should in principle be<br>aligned already. We can add padding in various structs, but according<br>to the spec adding float2 to differential3 shouldn't actually improve<br>alignment, so I'm not sure what the rule would be.<br><br><blockquote type="cite">Regarding (2) viewport problem:<br></blockquote><blockquote type="cite">----------------------------------------------<br></blockquote><blockquote type="cite">Does F12 rendering work (with resolutions that fail in the viewport) or is<br></blockquote><blockquote type="cite">it just viewport rendering? ...<br></blockquote><blockquote type="cite"><blockquote type="cite"> <a href="http://www.pasteall.org/27116/diff">http://www.pasteall.org/27116/diff</a><br></blockquote></blockquote><blockquote type="cite">The dumped image is as bad as the viewport one. A note: the image is not<br></blockquote><blockquote type="cite">'broken' right away. While the light calculation is still dirty, the image<br></blockquote><blockquote type="cite">is correct. Only after a few (set_tile() resolution > 8 or so) moments the<br></blockquote><blockquote type="cite">viewport breaks.<br></blockquote><br>Ok, so it's not an opengl issue.<br><br><blockquote type="cite"><blockquote type="cite">Another thing you could test is the workgroup_size in device_opencl.cpp in<br></blockquote></blockquote><blockquote type="cite"><blockquote type="cite">path_trace and tonemap.<br></blockquote></blockquote><blockquote type="cite">workgroup_size is calculated as 16. If I change it to 1, 2, 4, 8, 16 it<br></blockquote><blockquote type="cite">produces the same result.<br></blockquote><blockquote type="cite">If I change it to 32 OpenCL throw an error on me (OpenCL error (-54):<br></blockquote><blockquote type="cite">Invalid work group size)<br></blockquote><br>Ok, so workgroup size is probably not the issue.<br><br>Brecht.<br>_______________________________________________<br>Bf-cycles mailing list<br><a href="mailto:Bf-cycles@blender.org">Bf-cycles@blender.org</a><br>http://lists.blender.org/mailman/listinfo/bf-cycles<br></div></blockquote></div><br><div>
<span class="Apple-style-span" style="border-collapse: separate; color: rgb(0, 0, 0); font-family: Helvetica; font-style: normal; font-variant: normal; font-weight: normal; letter-spacing: normal; line-height: normal; orphans: 2; text-align: -webkit-auto; text-indent: 0px; text-transform: none; white-space: normal; widows: 2; word-spacing: 0px; -webkit-border-horizontal-spacing: 0px; -webkit-border-vertical-spacing: 0px; -webkit-text-decorations-in-effect: none; -webkit-text-size-adjust: auto; -webkit-text-stroke-width: 0px; font-size: medium; "><div><div><span class="Apple-style-span" style="font-size: 12px; ">_____________________________________</span></div><div><span class="Apple-style-span" style="font-size: 12px; "><div><br class="khtml-block-placeholder"></div><div style="text-align: left; ">Jens Verwiebe</div><div style="text-align: left; ">Allerskehre 44 - 22309 Hamburg</div><div style="text-align: left; "><br class="khtml-block-placeholder"></div><div style="text-align: left; ">Tel.: +49 40 68 78 50</div><div style="text-align: left; ">mobil: +49 172 400 49 07</div><div>mailto: <a href="mailto:info@jensverwiebe.de"><span class="Apple-style-span" style="color: rgb(0, 0, 238); ">info@jensverwiebe.de</span></a></div><div>web: <a href="http://www.jensverwiebe.de/"><span class="Apple-style-span" style="color: rgb(0, 0, 238); ">http://www.jensverwiebe.de</span></a></div><div style="text-align: left; ">_____________________________________</div></span></div></div></span>
</div>
<br></div></body></html>