<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 "define float3" patch is not applied. Otherwise I get an alpha viewport. Also when it'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>"Compiling OpenCL kernel ...</div><div><b>[a few seconds later ...]</b></div>Error getting function data from server"</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'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 "rip" the viewport apart the new editor only renders 'alpha'</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"><<a href="mailto:brechtvanlommel@pandora.be" target="_blank">brechtvanlommel@pandora.be</a>></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's based on doing blocks of max N pixels instead. That<br>
should make it work then. I couldn'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 <<a href="mailto:dfelinto@gmail.com" target="_blank">dfelinto@gmail.com</a>> wrote:<br>
> Hi Brecht,<br>
><br>
> Your latest patch [1] fix the problem only partially. Now it doesn't break<br>
> with 1280x1280 and lower, but in 2600x2600 it starts to crash. Funny fact,<br>
> with 2550x2550 or 2570x2570 I get a black bar on top [2].<br>
><br>
> Jens said that in Sig (lux and small lux) they are splitting the tasks as<br>
> well. So it's a good direction to go.<br>
> That still doesn't solve the problem of shaders not building as Jens<br>
> reported to me. I wonder if it's already time to address that.<br>
><br>
> Cheers,<br>
> Dalai<br>
><br>
><br>
> [1] - <a href="http://www.pasteall.org/27291/diff" target="_blank">http://www.pasteall.org/27291/diff</a><br>
> [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>
> 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>
><br>
><br>
> 2011/12/8 Dalai Felinto <<a href="mailto:dfelinto@gmail.com" target="_blank">dfelinto@gmail.com</a>><br>
>><br>
>> Hi,<br>
>> It seems that the viewport problem is related to ShaderData or<br>
>> ShaderClosure alignment or other problems.<br>
>> The problem is pretty much narrowed down to the following:<br>
>><br>
>> (1) in kernel_path.h : both lines 275 and 277 can break if not commented<br>
>> from the code:<br>
>> - shader_setup_from_ray(kg, &sd, &isect, &ray);<br>
>> - shader_eval_surface(kg, &sd, rbsdf, state.flag);<br>
>><br>
>> (2) in kernel_shader.h : shader_eval_surface() consist of only 2 lines of<br>
>> code (for opencl). Both lines (486 and 487) can also break if not commented<br>
>> out:<br>
>> - bsdf_diffuse_setup(sd, &sd->closure);<br>
>> - sd->closure.weight = make_float3(0.8f, 0.8f, 0.8f);<br>
>><br>
>> Notes:<br>
>> --------<br>
>> 1) by break I mean, "breaks" the viewport or get a pitch black render<br>
>> 2) resolutions used for testing(Width x Height @ Tiles): 820x772@820<br>
>> (breaks), 820x772@818 (works)<br>
>> 3) by "both lines need to be commented" I mean, if either one is left<br>
>> uncommented it breaks.<br>
>> 4) I get the same results with or without the kernel_compat_cl.h patch for<br>
>> #define float3 float4<br>
>> 5) my test file - <a href="http://www.pasteall.org/blend/10191" target="_blank">http://www.pasteall.org/blend/10191</a><br>
>> 6) example of "working" image when commenting the above mentioned lines:<br>
>> <a href="http://www.pasteall.org/pic/show.php?id=22245" target="_blank">http://www.pasteall.org/pic/show.php?id=22245</a><br>
>><br>
>> Best regards,<br>
>> Dalai<br>
>><br>
>> -- getting close ;)<br>
>><br>
>> 2011/12/7 Jens Verwiebe <<a href="mailto:info@jensverwiebe.de" target="_blank">info@jensverwiebe.de</a>><br>
>>><br>
>>> Hi<br>
>>> As a hint i found these 2 writings:<br>
>>><br>
>>> So about that float padding[2]; in the struct definition. This is because<br>
>>> of memoryalignment in OpenCL. The best explanation I’ve seen so far is<br>
>>> by AndreasStahl which I will briefly summarize in relation to the struct<br>
>>> above.<br>
>>> When interpreting a struct, OpenCL accesses the memory in blocks of 16<br>
>>> bytes, which is the same as 4 floats (each 4 bytes). So in our example if we<br>
>>> did not have the padding, we would not be able to access our int because<br>
>>> opencl would have interpreted it as the 3rd float out of the first 16bytes.<br>
>>> This can get even more complicated if you have an array of structs, because<br>
>>> then the size of you’re struct will need to be a multiple of 16, as<br>
>>> explained in the linked forum post.<br>
>>> The orig AMD posting:<br>
>>><br>
>>><br>
>>> <a href="http://forums.amd.com/forum/messageview.cfm?catid=390&threadid=122209&forumid=9" target="_blank">http://forums.amd.com/forum/messageview.cfm?catid=390&threadid=122209&forumid=9</a><br>
>>><br>
>>> Conclusion: There are indeed places where the compiler does not take care<br>
>>> of alignment ( bug? ), so we must do ourselfes.<br>
>>><br>
>>> Jens<br>
>>><br>
>>><br>
>>><br>
>>><br>
>>><br>
>>><br>
>>> Am 07.12.2011 um 16:38 schrieb Brecht Van Lommel:<br>
>>><br>
>>> 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" target="_blank">dfelinto@gmail.com</a>> wrote:<br>
>>><br>
>>> "Good news", I believe both problems may indeed be related and of<br>
>>> alignment.<br>
>>><br>
>>> At least the first one.<br>
>>><br>
>>><br>
>>> If I change the alignment in kernel_type.h in simple ways [*] I get the Z<br>
>>><br>
>>> problem fixed, but the viewport broken in situations it was working<br>
>>> before.<br>
>>><br>
>>> If I go to camera fly mode the viewport is fine ('dirty' but fine) until<br>
>>> I<br>
>>><br>
>>> stop then I get the (2) problem. So what are the rules for alignment<br>
>>> here?<br>
>>><br>
>>><br>
>>> [*] <a href="http://www.pasteall.org/27124" target="_blank">http://www.pasteall.org/27124</a><br>
>>><br>
>>><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>
>>> From the specification:<br>
>>><br>
>>><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>
>>> Regarding (2) viewport problem:<br>
>>><br>
>>> ----------------------------------------------<br>
>>><br>
>>> Does F12 rendering work (with resolutions that fail in the viewport) or<br>
>>> is<br>
>>><br>
>>> it just viewport rendering? ...<br>
>>><br>
>>> <a href="http://www.pasteall.org/27116/diff" target="_blank">http://www.pasteall.org/27116/diff</a><br>
>>><br>
>>> The dumped image is as bad as the viewport one. A note: the image is not<br>
>>><br>
>>> 'broken' right away. While the light calculation is still dirty, the<br>
>>> image<br>
>>><br>
>>> is correct. Only after a few (set_tile() resolution > 8 or so) moments<br>
>>> the<br>
>>><br>
>>> viewport breaks.<br>
>>><br>
>>><br>
>>> Ok, so it's not an opengl issue.<br>
>>><br>
>>> Another thing you could test is the workgroup_size in device_opencl.cpp<br>
>>> in<br>
>>><br>
>>> path_trace and tonemap.<br>
>>><br>
>>> workgroup_size is calculated as 16. If I change it to 1, 2, 4, 8, 16 it<br>
>>><br>
>>> produces the same result.<br>
>>><br>
>>> If I change it to 32 OpenCL throw an error on me (OpenCL error (-54):<br>
>>><br>
>>> Invalid work group size)<br>
>>><br>
>>><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" 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>
>>><br>
>>><br>
>>> _____________________________________<br>
>>><br>
>>> Jens Verwiebe<br>
>>> Allerskehre 44 - 22309 Hamburg<br>
>>><br>
>>> Tel.: <a href="tel:%2B49%2040%2068%2078%2050" value="+4940687850" target="_blank">+49 40 68 78 50</a><br>
>>> mobil: <a href="tel:%2B49%20172%20400%2049%2007" value="+491724004907" target="_blank">+49 172 400 49 07</a><br>
>>> mailto: <a href="mailto:info@jensverwiebe.de" target="_blank">info@jensverwiebe.de</a><br>
>>> web: <a href="http://www.jensverwiebe.de" target="_blank">http://www.jensverwiebe.de</a><br>
>>> _____________________________________<br>
>>><br>
>>><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>
>>><br>
>><br>
><br>
><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>
><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>