<div>Hi,</div><div>I have good and bad news ...</div><div><br></div><div>In my computer the kernel options that work are:</div><div>#define __SVM__</div><div>#define __TRANSPARENT_SHADOWS__</div><div><br></div><div>That only if the &quot;define float3&quot; patch is not applied. Otherwise I get an alpha viewport. Also when it&#39;s applied I have this warning, not sure is related:</div>

<div><a href="http://www.pasteall.org/27310">http://www.pasteall.org/27310</a></div><div><br></div><div>Regardless of this we have all this warnings going on:</div><div><a href="http://www.pasteall.org/27311">http://www.pasteall.org/27311</a></div>

<div><br></div><div>Note, enabling any of the other options (e.g. __EMISSION__) I get:</div><div><div>&quot;Compiling OpenCL kernel ...</div><div><b>[a few seconds later ...]</b></div>Error getting function data from server&quot;</div>

<div><br></div><div>Some options take a few seconds and throw this problem, others take more time.</div><div><br></div><div>The minimalistic patch I have with SVM working is:</div><div><a href="http://www.pasteall.org/27313/diff">http://www.pasteall.org/27313/diff</a></div>

<div><div><br></div><div>A screenshot to prove I&#39;m not lying:</div><div><a href="http://www.pasteall.org/pic/22421">http://www.pasteall.org/pic/22421</a></div><div>Now if I &quot;rip&quot; the viewport apart the new editor only renders &#39;alpha&#39;</div>

<div><a href="http://www.pasteall.org/pic/22422">http://www.pasteall.org/pic/22422</a></div><div><br></div><div>Blender test file:</div><div><a href="http://www.pasteall.org/blend/10223">http://www.pasteall.org/blend/10223</a></div>

<div><br></div><div>--</div><div>Dalai<br><br><div class="gmail_quote">2011/12/10 Brecht Van Lommel <span dir="ltr">&lt;<a href="mailto:brechtvanlommel@pandora.be" target="_blank">brechtvanlommel@pandora.be</a>&gt;</span><br>


<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Hi,<br>
<br>
I set it to split into a fixed number of parts now, but will add code<br>
so that it&#39;s based on doing blocks of max N pixels instead. That<br>
should make it work then. I couldn&#39;t find anything in the opencl spec<br>
about global range limitations for execution, but is a simple enough<br>
fix.<br>
<br>
What is exactly the problem of shaders not building, that it does only<br>
ambient occlusion render? In kernel_types.h, under /* Kernel Features<br>
*/, you can see which features are enabled. In device_opencl.cpp, it<br>
enables more for NVidia drivers. For Apple drivers, more features can<br>
be enabled too when they are verified to work.<br>
<br>
Thanks,<br>
<font color="#888888">Brecht.<br>
</font><div><div></div><div><br>
On Sat, Dec 10, 2011 at 9:39 AM, Dalai Felinto &lt;<a href="mailto:dfelinto@gmail.com" target="_blank">dfelinto@gmail.com</a>&gt; wrote:<br>
&gt; Hi Brecht,<br>
&gt;<br>
&gt; Your latest patch [1] fix the problem only partially. Now it doesn&#39;t break<br>
&gt; with 1280x1280 and lower, but in 2600x2600 it starts to crash. Funny fact,<br>
&gt; with 2550x2550 or 2570x2570 I get a black bar on top [2].<br>
&gt;<br>
&gt; Jens said that in Sig (lux and small lux) they are splitting the tasks as<br>
&gt; well. So it&#39;s a good direction to go.<br>
&gt; That still doesn&#39;t solve the problem of shaders not building as Jens<br>
&gt; reported to me. I wonder if it&#39;s already time to address that.<br>
&gt;<br>
&gt; Cheers,<br>
&gt; Dalai<br>
&gt;<br>
&gt;<br>
&gt; [1] - <a href="http://www.pasteall.org/27291/diff" target="_blank">http://www.pasteall.org/27291/diff</a><br>
&gt; [2] - <a href="http://www.pasteall.org/pic/show.php?id=22390" target="_blank">http://www.pasteall.org/pic/show.php?id=22390</a> in oppose<br>
&gt; to <a href="http://www.pasteall.org/pic/show.php?id=22391" target="_blank">http://www.pasteall.org/pic/show.php?id=22391</a><br>
&gt;<br>
&gt;<br>
&gt; 2011/12/8 Dalai Felinto &lt;<a href="mailto:dfelinto@gmail.com" target="_blank">dfelinto@gmail.com</a>&gt;<br>
&gt;&gt;<br>
&gt;&gt; Hi,<br>
&gt;&gt; It seems that the viewport problem is related to ShaderData or<br>
&gt;&gt; ShaderClosure alignment or other problems.<br>
&gt;&gt; The problem is pretty much narrowed down to the following:<br>
&gt;&gt;<br>
&gt;&gt; (1) in kernel_path.h : both lines 275 and 277 can break if not commented<br>
&gt;&gt; from the code:<br>
&gt;&gt; - shader_setup_from_ray(kg, &amp;sd, &amp;isect, &amp;ray);<br>
&gt;&gt; - shader_eval_surface(kg, &amp;sd, rbsdf, state.flag);<br>
&gt;&gt;<br>
&gt;&gt; (2) in kernel_shader.h : shader_eval_surface() consist of only 2 lines of<br>
&gt;&gt; code (for opencl). Both lines (486 and 487) can also break if not commented<br>
&gt;&gt; out:<br>
&gt;&gt;  - bsdf_diffuse_setup(sd, &amp;sd-&gt;closure);<br>
&gt;&gt; -  sd-&gt;closure.weight = make_float3(0.8f, 0.8f, 0.8f);<br>
&gt;&gt;<br>
&gt;&gt; Notes:<br>
&gt;&gt; --------<br>
&gt;&gt; 1) by break I mean, &quot;breaks&quot; the viewport or get a pitch black render<br>
&gt;&gt; 2) resolutions used for testing(Width x Height @ Tiles): 820x772@820<br>
&gt;&gt; (breaks), 820x772@818 (works)<br>
&gt;&gt; 3) by &quot;both lines need to be commented&quot; I mean, if either one is left<br>
&gt;&gt; uncommented it breaks.<br>
&gt;&gt; 4) I get the same results with or without the kernel_compat_cl.h patch for<br>
&gt;&gt; #define float3 float4<br>
&gt;&gt; 5) my test file - <a href="http://www.pasteall.org/blend/10191" target="_blank">http://www.pasteall.org/blend/10191</a><br>
&gt;&gt; 6) example of &quot;working&quot; image when commenting the above mentioned lines:<br>
&gt;&gt; <a href="http://www.pasteall.org/pic/show.php?id=22245" target="_blank">http://www.pasteall.org/pic/show.php?id=22245</a><br>
&gt;&gt;<br>
&gt;&gt; Best regards,<br>
&gt;&gt; Dalai<br>
&gt;&gt;<br>
&gt;&gt; -- getting close ;)<br>
&gt;&gt;<br>
&gt;&gt; 2011/12/7 Jens Verwiebe &lt;<a href="mailto:info@jensverwiebe.de" target="_blank">info@jensverwiebe.de</a>&gt;<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; Hi<br>
&gt;&gt;&gt; As a hint i found these 2 writings:<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; So about that float padding[2]; in the struct definition. This is because<br>
&gt;&gt;&gt; of memoryalignment in OpenCL. The best explanation I’ve seen so far is<br>
&gt;&gt;&gt; by AndreasStahl which I will briefly summarize in relation to the struct<br>
&gt;&gt;&gt; above.<br>
&gt;&gt;&gt; When interpreting a struct, OpenCL accesses the memory in blocks of 16<br>
&gt;&gt;&gt; bytes, which is the same as 4 floats (each 4 bytes). So in our example if we<br>
&gt;&gt;&gt; did not have the padding, we would not be able to access our int because<br>
&gt;&gt;&gt; opencl would have interpreted it as the 3rd float out of the first 16bytes.<br>
&gt;&gt;&gt; This can get even more complicated if you have an array of structs, because<br>
&gt;&gt;&gt; then the size of you’re struct will need to be a multiple of 16, as<br>
&gt;&gt;&gt; explained in the linked forum post.<br>
&gt;&gt;&gt; The orig AMD posting:<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; <a href="http://forums.amd.com/forum/messageview.cfm?catid=390&amp;threadid=122209&amp;forumid=9" target="_blank">http://forums.amd.com/forum/messageview.cfm?catid=390&amp;threadid=122209&amp;forumid=9</a><br>



&gt;&gt;&gt;<br>
&gt;&gt;&gt; Conclusion: There are indeed places where the compiler does not take care<br>
&gt;&gt;&gt; of alignment ( bug? ), so we must do ourselfes.<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; Jens<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; Am 07.12.2011 um 16:38 schrieb Brecht Van Lommel:<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; Hi,<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; Maybe it would be possible to set up an SSH login on a computer so I<br>
&gt;&gt;&gt; can try and narrow this down further? There has to be a pattern here,<br>
&gt;&gt;&gt; but I&#39;m not sure what to suggest you to test, what I would do is to<br>
&gt;&gt;&gt; keep removing code until it works, and find out exactly which kind of<br>
&gt;&gt;&gt; struct variable layout or operation is causing issues.<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; On Wed, Dec 7, 2011 at 7:11 AM, Dalai Felinto &lt;<a href="mailto:dfelinto@gmail.com" target="_blank">dfelinto@gmail.com</a>&gt; wrote:<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; &quot;Good news&quot;, I believe both problems may indeed be related and of<br>
&gt;&gt;&gt; alignment.<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; At least the first one.<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; If I change the alignment in kernel_type.h in simple ways [*] I get the Z<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; problem fixed, but the viewport broken in situations it was working<br>
&gt;&gt;&gt; before.<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; If I go to camera fly mode the viewport is fine (&#39;dirty&#39; but fine) until<br>
&gt;&gt;&gt; I<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; stop then I get the (2) problem. So what are the rules for alignment<br>
&gt;&gt;&gt; here?<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; [*] <a href="http://www.pasteall.org/27124" target="_blank">http://www.pasteall.org/27124</a><br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; The alignment issue I was thinking of is when passing structs between<br>
&gt;&gt;&gt; the C++ and OpenCL. If they layout variables in structs differently,<br>
&gt;&gt;&gt; copying KernelData to the GPU would go wrong. So what I&#39;ve tried to do<br>
&gt;&gt;&gt; is align everything like we do in makesdna, and make all struct sizes<br>
&gt;&gt;&gt; a multiple of 16 bytes.<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; What is quite strange, is that the struct where you added alignment<br>
&gt;&gt;&gt; does not get passed between C++ and OpenCL, it stays purely in OpenCL,<br>
&gt;&gt;&gt; and so I wouldn&#39;t expect this to be a problem. Maybe there&#39;s a<br>
&gt;&gt;&gt; compiler bug related to nested structs?<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; From the specification:<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; &quot;The OpenCL compiler is responsible for aligning data items to the<br>
&gt;&gt;&gt; appropriate alignment as required by the data type. For arguments to a<br>
&gt;&gt;&gt; __kernel function declared to be a pointer to a data type, the OpenCL<br>
&gt;&gt;&gt; compiler can assume that the pointee is always appropriately aligned<br>
&gt;&gt;&gt; as required by the data type.&quot;<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; So this means that in principle we shouldn&#39;t have to worry about<br>
&gt;&gt;&gt; alignment of structs that stay on the GPU, and only for KernelData we<br>
&gt;&gt;&gt; need to worry about alignment. But of course compiler bugs can make<br>
&gt;&gt;&gt; this an issue.<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; &quot;For 3-component vector data types, the size of the data type is 4 *<br>
&gt;&gt;&gt; sizeof(component). This means that a 3-component vector data type will<br>
&gt;&gt;&gt; be aligned to a 4 * sizeof(component) boundary.&quot;<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; This seems to indicate that differential3 should in principle be<br>
&gt;&gt;&gt; aligned already. We can add padding in various structs, but according<br>
&gt;&gt;&gt; to the spec adding float2 to differential3 shouldn&#39;t actually improve<br>
&gt;&gt;&gt; alignment, so I&#39;m not sure what the rule would be.<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; Regarding (2) viewport problem:<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; ----------------------------------------------<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; Does F12 rendering work (with resolutions that fail in the viewport) or<br>
&gt;&gt;&gt; is<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; it just viewport rendering? ...<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt;  <a href="http://www.pasteall.org/27116/diff" target="_blank">http://www.pasteall.org/27116/diff</a><br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; The dumped image is as bad as the viewport one. A note: the image is not<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; &#39;broken&#39; right away. While the light calculation is still dirty, the<br>
&gt;&gt;&gt; image<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; is correct. Only after a few (set_tile() resolution &gt; 8 or so) moments<br>
&gt;&gt;&gt; the<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; viewport breaks.<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; Ok, so it&#39;s not an opengl issue.<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; Another thing you could test is the workgroup_size in device_opencl.cpp<br>
&gt;&gt;&gt; in<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; path_trace and tonemap.<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; workgroup_size is calculated as 16. If I change it to 1, 2, 4, 8, 16 it<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; produces the same result.<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; If I change it to 32 OpenCL throw an error on me (OpenCL error (-54):<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; Invalid work group size)<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; Ok, so workgroup size is probably not the issue.<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; Brecht.<br>
&gt;&gt;&gt; _______________________________________________<br>
&gt;&gt;&gt; Bf-cycles mailing list<br>
&gt;&gt;&gt; <a href="mailto:Bf-cycles@blender.org" target="_blank">Bf-cycles@blender.org</a><br>
&gt;&gt;&gt; <a href="http://lists.blender.org/mailman/listinfo/bf-cycles" target="_blank">http://lists.blender.org/mailman/listinfo/bf-cycles</a><br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; _____________________________________<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; Jens Verwiebe<br>
&gt;&gt;&gt; Allerskehre 44  -  22309 Hamburg<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; Tel.: <a href="tel:%2B49%2040%2068%2078%2050" value="+4940687850" target="_blank">+49 40 68 78 50</a><br>
&gt;&gt;&gt; mobil: <a href="tel:%2B49%20172%20400%2049%2007" value="+491724004907" target="_blank">+49 172 400 49 07</a><br>
&gt;&gt;&gt; mailto: <a href="mailto:info@jensverwiebe.de" target="_blank">info@jensverwiebe.de</a><br>
&gt;&gt;&gt; web:  <a href="http://www.jensverwiebe.de" target="_blank">http://www.jensverwiebe.de</a><br>
&gt;&gt;&gt; _____________________________________<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt;<br>
&gt;&gt;&gt; _______________________________________________<br>
&gt;&gt;&gt; Bf-cycles mailing list<br>
&gt;&gt;&gt; <a href="mailto:Bf-cycles@blender.org" target="_blank">Bf-cycles@blender.org</a><br>
&gt;&gt;&gt; <a href="http://lists.blender.org/mailman/listinfo/bf-cycles" target="_blank">http://lists.blender.org/mailman/listinfo/bf-cycles</a><br>
&gt;&gt;&gt;<br>
&gt;&gt;<br>
&gt;<br>
&gt;<br>
&gt; _______________________________________________<br>
&gt; Bf-cycles mailing list<br>
&gt; <a href="mailto:Bf-cycles@blender.org" target="_blank">Bf-cycles@blender.org</a><br>
&gt; <a href="http://lists.blender.org/mailman/listinfo/bf-cycles" target="_blank">http://lists.blender.org/mailman/listinfo/bf-cycles</a><br>
&gt;<br>
_______________________________________________<br>
Bf-cycles mailing list<br>
<a href="mailto:Bf-cycles@blender.org" target="_blank">Bf-cycles@blender.org</a><br>
<a href="http://lists.blender.org/mailman/listinfo/bf-cycles" target="_blank">http://lists.blender.org/mailman/listinfo/bf-cycles</a><br>
</div></div></blockquote></div><br></div></div>