ATI released OpenCL SDK with hardware support

[jensverwiebe]

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

Here´s the result:

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

Jens

[Dade]

@ 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.

[jensverwiebe]

@ 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

[Dade]

@ 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

[jensverwiebe]

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

[Dade]

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.

[Dade]

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):

[jensverwiebe]

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..


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

[Eros]

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.

[Dade]

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";