LuxCore: new BIASPATHOCL with OpenCL micro-kernels

Discussion related to the implementation of new features & algorithms to the Core Engine.

Moderators: Dade, jromang, tomb, zcott, coordinators

Re: LuxCore: new BIASPATHOCL with OpenCL micro-kernels

Postby jensverwiebe » Mon Oct 10, 2016 10:01 am

Dade wrote:
jensverwiebe wrote:Edit, found something new in message when going with luxcoreui:
Code: Select all
[LuxRays][6.730] [Device GeForce GTX 980 Intersect] RADIANCE_PER_PIXEL_NORMALIZEDs[0] buffer size: 4800Kbytes
[LuxCore][6.730] [PathOCLRenderThread::0] Rendering thread ERROR: clEnqueueNDRangeKernel(CL_OUT_OF_RESOURCES)


Thats the culprit ?


CL_OUT_OF_RESOURCES is an error NVIDIA driver usually throws when ... it doesn't know what is going wrong on the GPU. It is a bit like a memory fault exception with the CPU.

Anyway I should have fixed the problem with my last commit and now Metropolis should works fine again.


Yo man, checked all okay now, ty :D
( taking OUT_OF_RESOURCES in its strict meaning wouldn't have made any sense here anyway :lol: )

Jens
User avatar
jensverwiebe
Developer
 
Posts: 3429
Joined: Wed Apr 02, 2008 4:34 pm

Re: LuxCore: new BIASPATHOCL with OpenCL micro-kernels

Postby jensverwiebe » Mon Oct 10, 2016 3:11 pm

Okay, exept preview in viewport ( which is not urgent ) all looks fine.
Trying to render the blender viewport i get new warnings and 2 errors i just wanna list here to work on later:

Code: Select all
[LuxCore][57.277] [PathOCLBaseRenderThread::0] Compiling kernels
[LuxCore][57.393] [PathOCLBaseRenderThread::0] PathOCL kernel compilation error
ERROR clBuildProgram[CL_BUILD_PROGRAM_FAILURE]:
<kernel>:7973:96: warning: passing 'float const __global[16]' to parameter of type '__global float *' discards qualifiers
        return BandTexture_ConstEvaluateFloat(hitPoint, texture->band.interpType, texture->band.size, texture->band.offsets, texture->band.values, ImageMapTexture_ConstEvaluateFloat(&texs[6], hitPoint IMAGEMAPS_PARAM));
                                                                                                      ^~~~~~~~~~~~~~~~~~~~~
<kernel>:7870:36: note: passing argument to parameter 'offsets' here
                const uint size, __global float *offsets,
                                                 ^
<kernel>:7973:119: warning: passing 'Spectrum const __global[16]' to parameter of type '__global Spectrum *' (aka '__global RGBColor *') discards qualifiers
        return BandTexture_ConstEvaluateFloat(hitPoint, texture->band.interpType, texture->band.size, texture->band.offsets, texture->band.values, ImageMapTexture_ConstEvaluateFloat(&texs[6], hitPoint IMAGEMAPS_PARAM));
                                                                                                                             ^~~~~~~~~~~~~~~~~~~~
<kernel>:7871:22: note: passing argument to parameter 'values' here
                __global Spectrum *values, const float amt) {
                                   ^
<kernel>:7978:99: warning: passing 'float const __global[16]' to parameter of type '__global float *' discards qualifiers
        return BandTexture_ConstEvaluateSpectrum(hitPoint, texture->band.interpType, texture->band.size, texture->band.offsets, texture->band.values, ImageMapTexture_ConstEvaluateFloat(&texs[6], hitPoint IMAGEMAPS_PARAM));
                                                                                                         ^~~~~~~~~~~~~~~~~~~~~
<kernel>:7829:36: note: passing argument to parameter 'offsets' here
                const uint size, __global float *offsets,
                                                 ^
<kernel>:7978:122: warning: passing 'Spectrum const __global[16]' to parameter of type '__global Spectrum *' (aka '__global RGBColor *') discards qualifiers
        return BandTexture_ConstEvaluateSpectrum(hitPoint, texture->band.interpType, texture->band.size, texture->band.offsets, texture->band.values, ImageMapTexture_ConstEvaluateFloat(&texs[6], hitPoint IMAGEMAPS_PARAM));
                                                                                                                                ^~~~~~~~~~~~~~~~~~~~
<kernel>:7830:22: note: passing argument to parameter 'values' here
                __global Spectrum *values, const float3 amt) {
                                   ^
<kernel>:16744:36: error: no member named 'pixelX' in 'SampleResult'
                                const uint px = sample->result.pixelX + x;
                                                ~~~~~~~~~~~~~~ ^
<kernel>:16745:36: error: no member named 'pixelY' in 'SampleResult'
                                const uint py = sample->result.pixelY + y;
                                                ~~~~~~~~~~~~~~ ^
<kernel>:16746:13: warning: comparison of unsigned expression >= 0 is always true
                                if ((px >= 0) && (px < filmWidth) && (py >= 0) && (py < filmHeight)) {
                                     ~~ ^  ~
<kernel>:16746:46: warning: comparison of unsigned expression >= 0 is always true
                                if ((px >= 0) && (px < filmWidth) && (py >= 0) && (py < filmHeight)) {
                                                                      ~~ ^  ~
<kernel>:16754:33: error: no member named 'pixelX' in 'SampleResult'
                Film_AddSample(sample->result.pixelX, sample->result.pixelY,
                               ~~~~~~~~~~~~~~ ^
<kernel>:17792:32: warning: use of logical '&&' with constant operand
        return ((component & DIFFUSE) && (PARAM_MAX_PATH_DEPTH_DIFFUSE > 0)) ||
                                      ^  ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
<kernel>:17792:32: note: use '&' for a bitwise operation
        return ((component & DIFFUSE) && (PARAM_MAX_PATH_DEPTH_DIFFUSE > 0)) ||
                                      ^~
                                      &
<kernel>:17792:32: note: remove constant to silence this warning
        return ((component & DIFFUSE) && (PARAM_MAX_PATH_DEPTH_DIFFUSE > 0)) ||
                                     ~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
<kernel>:17793:26: warning: use of logical '&&' with constant operand
                        ((component & GLOSSY) && (PARAM_MAX_PATH_DEPTH_GLOSSY > 0)) ||
                                              ^  ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
<kernel>:17793:26: note: use '&' for a bitwise operation
                        ((component & GLOSSY) && (PARAM_MAX_PATH_DEPTH_GLOSSY > 0)) ||
                                              ^~
                                              &
<kernel>:17793:26: note: remove constant to silence this warning
                        ((component & GLOSSY) && (PARAM_MAX_PATH_DEPTH_GLOSSY > 0)) ||
                                             ~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
<kernel>:17794:28: warning: use of logical '&&' with constant operand
                        ((component & SPECULAR) && (PARAM_MAX_PATH_DEPTH_SPECULAR > 0));
                                                ^  ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
<kernel>:17794:28: note: use '&' for a bitwise operation
                        ((component & SPECULAR) && (PARAM_MAX_PATH_DEPTH_SPECULAR > 0));
                                                ^~
                                                &
<kernel>:17794:28: note: remove constant to silence this warning
                        ((component & SPECULAR) && (PARAM_MAX_PATH_DEPTH_SPECULAR > 0));
                                               ~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
<kernel>:17889:9: warning: type specifier missing, defaults to 'int'
                const pixelIndex = tilePass % pixelsCount2;
                ~~~~~ ^


Traceback (most recent call last):
  File "/home/jensverwiebe/.config/blender/2.77/scripts/addons/luxrender/core/__init__.py", line 446, in view_update
    self.luxcore_view_update(context)
  File "/home/jensverwiebe/.config/blender/2.77/scripts/addons/luxrender/core/__init__.py", line 2389, in luxcore_view_update
    LuxCoreSessionManager.start_luxcore_session(self.space)
  File "/home/jensverwiebe/.config/blender/2.77/scripts/addons/luxrender/core/__init__.py", line 2520, in start_luxcore_session
    session.luxcore_session.Start()
RuntimeError: PathOCLBase kernel compilation error

location: <unknown location>:-1

location: <unknown location>:-1


Not yet investigated which is caused in core or export. Again:only in viewport render !

EDIT: fixed the "bitwise" op complaints in ocl. Hope this goes in the right direction.
Down to ( occures only in viewport render ):
Code: Select all
[LuxCore][8.182] [PathOCLBaseRenderThread::0] Compiling kernels
[LuxCore][8.305] [PathOCLBaseRenderThread::0] PathOCL kernel compilation error
ERROR clBuildProgram[CL_BUILD_PROGRAM_FAILURE]:
<kernel>:16767:36: error: no member named 'pixelX' in 'SampleResult'
                                const uint px = sample->result.pixelX + x;
                                                ~~~~~~~~~~~~~~ ^
<kernel>:16768:36: error: no member named 'pixelY' in 'SampleResult'
                                const uint py = sample->result.pixelY + y;
                                                ~~~~~~~~~~~~~~ ^
<kernel>:16769:13: warning: comparison of unsigned expression >= 0 is always true
                                if ((px >= 0) && (px < filmWidth) && (py >= 0) && (py < filmHeight)) {
                                     ~~ ^  ~
<kernel>:16769:46: warning: comparison of unsigned expression >= 0 is always true
                                if ((px >= 0) && (px < filmWidth) && (py >= 0) && (py < filmHeight)) {
                                                                      ~~ ^  ~
<kernel>:16777:33: error: no member named 'pixelX' in 'SampleResult'
                Film_AddSample(sample->result.pixelX, sample->result.pixelY,
                               ~~~~~~~~~~~~~~ ^
<kernel>:17912:9: warning: type specifier missing, defaults to 'int'
                const pixelIndex = tilePass % pixelsCount2;
                ~~~~~ ^


[Lux 2016-Oct-11 14:37:11] View update aborted: PathOCLBase kernel compilation error


EDIT2: fixed also type specifier warning
Jens
User avatar
jensverwiebe
Developer
 
Posts: 3429
Joined: Wed Apr 02, 2008 4:34 pm

Re: LuxCore: new BIASPATHOCL with OpenCL micro-kernels

Postby Dade » Tue Oct 11, 2016 8:31 am

I should have fixed this problem, it was caused by using RTBIASPATHOCL with NONE filter (it was working fine with other pixel filters).

jensverwiebe wrote:EDIT: fixed the "bitwise" op complaints in ocl.


Nope, they are logical operators, I reverted the patch.
User avatar
Dade
Developer
 
Posts: 8404
Joined: Sat Apr 19, 2008 6:04 pm
Location: Italy

Re: LuxCore: new BIASPATHOCL with OpenCL micro-kernels

Postby jensverwiebe » Tue Oct 11, 2016 8:36 am

Ah, thx, again a misleading message. Had a bad feeling with that anyway, although it silenced the warnings :?:

EDIT: checked viewport render works now okay. Man these ocl logs are really misleading.
Kernel compile does not exhibit any warnings anymore.

At some point i did not lonnger had any idea where to look for :evil:
Hope ppl learn from my mistakes :oops:

@Dade: does the AMD sdk also show such misleading debug info ?


Jens
User avatar
jensverwiebe
Developer
 
Posts: 3429
Joined: Wed Apr 02, 2008 4:34 pm

Re: LuxCore: new BIASPATHOCL with OpenCL micro-kernels

Postby B.Y.O.B. » Tue Oct 11, 2016 9:05 am

Great work Dade, looking forward to try it out.
I'm also going to write a new thread with a proposal of "Path engine unification for exporters": showing a single PATH rendering engine with just several flags: CPU/OpenCL, Tile/No-tile, RT/No-RT, etc.

This also sounds very interesting.
User avatar
B.Y.O.B.
Developer
 
Posts: 5181
Joined: Wed Nov 10, 2010 4:10 pm
Location: Germany

Re: LuxCore: new BIASPATHOCL with OpenCL micro-kernels

Postby Dade » Tue Oct 11, 2016 1:50 pm

jensverwiebe wrote:@Dade: does the AMD sdk also show such misleading debug info ?


It fires few warnings, indeed, they are different from NVIDIA one and all wrong too :roll:
User avatar
Dade
Developer
 
Posts: 8404
Joined: Sat Apr 19, 2008 6:04 pm
Location: Italy

Re: LuxCore: new BIASPATHOCL with OpenCL micro-kernels

Postby kalel » Sat Oct 15, 2016 9:56 pm

Great work Dade, this brings some great stuff, also for PathOCL. I've updated my windows builds thread to include the current commits.

Is it possible also to not render diffuse, specular or glossy indirect light for specific materials/objects with Path, or does that remain a BiasPath exclusive feature?
kalel
 
Posts: 258
Joined: Sun Jul 17, 2016 6:35 am

Re: LuxCore: new BIASPATHOCL with OpenCL micro-kernels

Postby B.Y.O.B. » Sun Oct 16, 2016 1:13 am

kalel wrote:I've updated my windows builds thread to include the current commits.

Note that LuxBlend is not compatible with these latest changes yet if you use Biased Path (because it now requires a special sampler).
If you want to use Biased Path, you will have to add this to the custom properties in LuxBlend:
Code: Select all
sampler.type = BIASPATHSAMPLER
User avatar
B.Y.O.B.
Developer
 
Posts: 5181
Joined: Wed Nov 10, 2010 4:10 pm
Location: Germany

Re: LuxCore: new BIASPATHOCL with OpenCL micro-kernels

Postby kalel » Sun Oct 16, 2016 2:11 am

B.Y.O.B. wrote:
kalel wrote:I've updated my windows builds thread to include the current commits.

Note that LuxBlend is not compatible with these latest changes yet if you use Biased Path (because it now requires a special sampler).
If you want to use Biased Path, you will have to add this to the custom properties in LuxBlend:
Code: Select all
sampler.type = BIASPATHSAMPLER


Really? I don't mean about the new sampler being needed, I'm aware of that, but about Luxblend not supporting it yet. After hitting "update" on Luxblend, I'm not receiving that error anymore.
If you get some spare time, please confirm this with my build, because it works properly on my machine (so maybe my test was faulty for some reason, I don't know which, but mistakes are always possible).
If it doesn't work, I'll try to modify Luxblend temporarily, or just add that notice until you're able to.

I don't think the luxblend update can regress my build to non-split-kernel version, as pylux and pyluxcore "dll's'" (pyd's) remain the same. Hopefully something like that can't happen.

There's always possibility of other strange mistakes that I might have made on my side, but hopefully not.

Edit: Here's a small video of rendering with my build using biaspath ocl:
https://cdn-e1.streamable.com/video/mp4 ... 5d18aa72b3

Maybe from the console, you can conclude whether the microkernel version is used properly or not. I think that it is, but I'm not sure about what happens with the sampler especially if it's not supported.
Last edited by kalel on Sun Oct 16, 2016 2:32 am, edited 2 times in total.
kalel
 
Posts: 258
Joined: Sun Jul 17, 2016 6:35 am

Re: LuxCore: new BIASPATHOCL with OpenCL micro-kernels

Postby B.Y.O.B. » Sun Oct 16, 2016 2:29 am

You're right, I was not aware that Jens already implemented it: https://bitbucket.org/luxrender/luxblen ... 6d272a0073
A LuxBlend update will not touch your Lux binaries in any way, don't worry.
User avatar
B.Y.O.B.
Developer
 
Posts: 5181
Joined: Wed Nov 10, 2010 4:10 pm
Location: Germany

PreviousNext

Return to Architecture & Design

Who is online

Users browsing this forum: No registered users and 1 guest