[Bf-cycles] osx + ati + cycles, looking for insights in debugging

Fri Dec 9 07:02:27 CET 2011

Hi,
It seems that the viewport problem is related to ShaderData or
ShaderClosure alignment or other problems.
The problem is pretty much narrowed down to the following:

(1) in kernel_path.h : both lines 275 and 277 can break if not commented
from the code:
- shader_setup_from_ray(kg, &sd, &isect, &ray);
- shader_eval_surface(kg, &sd, rbsdf, state.flag);

(2) in kernel_shader.h : shader_eval_surface() consist of only 2 lines of
code (for opencl). Both lines (486 and 487) can also break if not commented
out:
 - bsdf_diffuse_setup(sd, &sd->closure);
-  sd->closure.weight = make_float3(0.8f, 0.8f, 0.8f);

Notes:
--------
1) by break I mean, "breaks" the viewport or get a pitch black render
2) resolutions used for testing(Width x Height @ Tiles):
820x772 at 820(breaks), 820x772 at 818(works)
3) by "both lines need to be commented" I mean, if either one is left
uncommented it breaks.
4) I get the same results with or without the kernel_compat_cl.h patch for
#define float3 float4
5) my test file - http://www.pasteall.org/blend/10191
6) example of "working" image when commenting the above mentioned lines:
http://www.pasteall.org/pic/show.php?id=22245

Best regards,
Dalai

-- getting close ;)

2011/12/7 Jens Verwiebe <info at jensverwiebe.de>

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