ATI released OpenCL SDK with hardware support

Discussions related to GPU Acceleration in LuxRender

Moderators: Dade, jromang, tomb, coordinators

Re: ATI released OpenCL SDK with hardware support

Postby jensverwiebe » Thu Jan 07, 2010 10:21 am

@ Dade...thx for helping out....

Here´s the result:

testkernel.jpg
./smallptGPU 1 32 preprocessed_rendering_kernel.cl 640 480 scenes/simple.scn



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

Re: ATI released OpenCL SDK with hardware support

Postby Dade » Thu Jan 07, 2010 11:00 am

jensverwiebe wrote:@ Dade...thx for helping out....

Here´s the result:
testkernel.jpg



Good, it is exactly how it should look. Now, let's check if some more buffer works with the following kernel:

Code: Select all
#include "camera.h"
#include "geomfunc.h"

__kernel void RadianceGPU(
    __global Vec *colors, __global unsigned int *seedsInput,
   __constant Sphere *sphere, __constant Camera *camera,
   const unsigned int sphereCount,
   const int width, const int height,
   const int currentSample,
   __global int *pixels,
   const unsigned int renderingFlags) {
    const int gid = get_global_id(0);
   const int x = gid % width;
   const int y = gid / width;

   /* Check if we have to do something */
   if (y >= height)
      return;

   Vec r;
   vinit(r, 0.001f, 0.001f, 0.001f);

   const int i = (height - y - 1) * width + x;
   if (currentSample == 0) {
      // Jens's patch for MacOS
      vinit(colors[i], 0.f, 0.f, 0.f);
   } else {
      vadd(colors[i], colors[i], r);

      if (colors[i].x > 1.f) {
         vclr(colors[i]);
      }
   }

   pixels[y * width + x] = toInt(colors[i].x) |
         (toInt(colors[i].y) << 8) |
         (toInt(colors[i].z) << 16);
}


This should produce a screen going from black to wait.
User avatar
Dade
Developer
 
Posts: 8318
Joined: Sat Apr 19, 2008 6:04 pm
Location: Italy

Re: ATI released OpenCL SDK with hardware support

Postby jensverwiebe » Thu Jan 07, 2010 11:06 am

@ Dade

With this kernel the render started uniform black and got brighter and brighter until it was pure white, then resetting to black and start over.....
i compared both cpu and gpu-mode, same behaviour.
Right this way?

BTW:
/* LordCRC: move seed to local store */
unsigned int seed0 = seedsInput[gid2];
unsigned int seed1 = seedsInput[gid2 + 1];


Could it be the missing local mem on older cards is the culprit?

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

Re: ATI released OpenCL SDK with hardware support

Postby Dade » Thu Jan 07, 2010 11:19 am

jensverwiebe wrote:@ Dade

With this kernel the render started uniform black and got brighter and brighter until it was pure white, then resetting to black and start over.....
i compared both cpu and gpu-mode, same behaviour.
Right this way?


Perfect, it looks like the color buffer and the pixel buffer are fine.

Jens, if you go back to the 1.6alpha and you try to resize the window, what does happen ? ... it could start to work :D
User avatar
Dade
Developer
 
Posts: 8318
Joined: Sat Apr 19, 2008 6:04 pm
Location: Italy

Re: ATI released OpenCL SDK with hardware support

Postby jensverwiebe » Thu Jan 07, 2010 11:26 am

Jens, if you go back to the 1.6alpha and you try to resize the window, what does happen ? ... it could start to work


Dade my Dade...unfortunately not.
I started with different initial sizes and resized with draging the lower corner, but only got some patterned distortions in some resizings, no correct renderoutput. :(

Keep in mind i got with cpu 1500k for cornell and 9000k for simple, but 75000k for cornell and simple in gpu mode, sounds not realistic to me, cornell should do around 12000k, no ?

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

Re: ATI released OpenCL SDK with hardware support

Postby Dade » Thu Jan 07, 2010 11:37 am

jensverwiebe wrote:
Jens, if you go back to the 1.6alpha and you try to resize the window, what does happen ? ... it could start to work


Dade my Dade...unfortunately not.
I started with different initial sizesand rezided, but only got some patterned distortions in some resizings, no correct renderoutput. :(


Ok, can you try the sources attached to this message ? They fill the initial pixel buffer with the usual colour pattern you have seen above. If the pattern stay there it means it is never overwritten.

P.S. you can try to resize too, it is an operation that reinitializes the buffers.
Attachments
smallptgpu-v1.6alpha.tgz
(157.52 KiB) Downloaded 50 times
User avatar
Dade
Developer
 
Posts: 8318
Joined: Sat Apr 19, 2008 6:04 pm
Location: Italy

Re: ATI released OpenCL SDK with hardware support

Postby Dade » Thu Jan 07, 2010 11:42 am

jensverwiebe wrote:Keep in mind i got with cpu 1500k for cornell and 9000k for simple, but 75000k for cornell and simple in gpu mode, sounds not realistic to me, cornell should do around 12000k, no ?


If the camera buffer is not downloaded to the GPU (it is the hypothesis I have at the moment), you could just be in the middle of nowhere looking to nothing. People are reaching 500,000,000 sample/sec looking to empty space (i.e. turning the camera in simple.scn in order to look at empty space):

Image
User avatar
Dade
Developer
 
Posts: 8318
Joined: Sat Apr 19, 2008 6:04 pm
Location: Italy

Re: ATI released OpenCL SDK with hardware support

Postby jensverwiebe » Thu Jan 07, 2010 11:45 am

k...compiled it and it behaves as you described:

At start the pattern shows up for a short time and disappers, but shows up again on every resize.

and....
you could just be in the middle of nowhere looking to nothing
.... please not now, we can do that later, much much later.. :o :shock: :roll: :lol:


hmm, could this have todo with preprocessing kernel ? do i need CPP flags perhaps ( like -fno-pic etc. ) ? This is a difference in our setups that comes to my eye atm.

Jens
Last edited by jensverwiebe on Thu Jan 07, 2010 12:37 pm, edited 1 time in total.
User avatar
jensverwiebe
Developer
 
Posts: 3402
Joined: Wed Apr 02, 2008 4:34 pm

Re: ATI released OpenCL SDK with hardware support

Postby Eros » Thu Jan 07, 2010 12:20 pm

I had a look at the first smallpt example posted earlier - the 1.6alpha - I see the same issue, only you said being in the middle of nowhere... when i use the CPU on OSX, the scenes render fine, on the GPU it is black. However in the simple scene I get a white dot kind of near centre of the window. I can move the camera around, which does move the dot as though it exists in 3D, trying to move towards it doesn't make it increase in size. It appears in all scenes, in the same place (i think, or at least roughly the same place.)

Not sure if knowing that helps at all - or is just another odd artefact.
User avatar
Eros
 
Posts: 418
Joined: Wed Jul 22, 2009 8:37 am

Re: ATI released OpenCL SDK with hardware support

Postby Dade » Thu Jan 07, 2010 1:10 pm

jensverwiebe wrote:At start the pattern shows up for a short time and disappers, but shows up again on every resize.


Ok, another kernel to check if the spheres array and camera are downloaded correctly:

Code: Select all
#define GPU_KERNEL

#include "camera.h"
#include "geomfunc.h"

__kernel void RadianceGPU(
    __global Vec *colors, __global unsigned int *seedsInput,
   __constant Sphere *sphere, __constant Camera *camera,
   const unsigned int sphereCount,
   const int width, const int height,
   const int currentSample,
   __global int *pixels,
   const unsigned int renderingFlags) {
    const int gid = get_global_id(0);
   const int x = gid % width;
   const int y = gid / width;

   /* Check if we have to do something */
   if (y >= height)
      return;

   int cameraOK = 1;
   if ((camera->orig.x != 50.f) ||
         (camera->orig.y != 45.f) ||
         (camera->orig.z != 205.6f) ||
         (camera->target.x != 50.f) ||
         (camera->target.y != 44.957388f) ||
         (camera->target.z != 204.6f))
      cameraOK = 0;

   int sceneOK = 1;
   if ((sphere[0].rad != 10000.f) || (sphere[1].rad != 10000.f) ||
         (sphere[0].p.x != 10001.f))
      sceneOK = 0;

   int col = 0;
   if (cameraOK && sceneOK)
      col = 0x00ff00;
   else if (!cameraOK && sceneOK)
      col = 0xff0000;
   else if (cameraOK && !sceneOK)
      col = 0x0000ff;
   else if (!cameraOK && !sceneOK)
      col = 0xff00ff;

   pixels[y * width + x] = col;
}


If you ran it you should the following screen colours:

green: all ok;
blue: camera wrong, scene ok
red: camera ok, scene wrong
pink: both wrong

if you move the camera (i.e. pressing a cursor key) or a sphere (i.e. pressing key '4') the related buffer should become "wrong";
User avatar
Dade
Developer
 
Posts: 8318
Joined: Sat Apr 19, 2008 6:04 pm
Location: Italy

PreviousNext

Return to GPU Acceleration

Who is online

Users browsing this forum: No registered users and 3 guests