RSS/Atom feed Twitter
Site is read-only, email is disabled

[PATCH 1/1] cl: Make noise-simplex kernel OpenCL 1.1 complaint

This discussion is connected to the gegl-developer-list.gnome.org mailing list which is provided by the GIMP developers and not related to gimpusers.com.

This is a read-only list on gimpusers.com so this discussion thread is read-only, too.

Jan Vesely
2014-05-18 20:38:08 UTC (almost 10 years ago)

[PATCH 1/1] cl: Make noise-simplex kernel OpenCL 1.1 complaint

No static keyword, f suffix for float constants Fixes one error and 7 warnings

Signed-off-by: Jan Vesely ---
Makes the op run on mesa/clover.

opencl/noise-simplex.cl | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-)

diff --git a/opencl/noise-simplex.cl b/opencl/noise-simplex.cl index bb59758..ce8aa40 100644
--- a/opencl/noise-simplex.cl
+++ b/opencl/noise-simplex.cl
@@ -1,6 +1,6 @@
#define MAX_RANK 3

-static float2
+float2
philox (uint2 st,
uint k)
{
@@ -17,7 +17,7 @@ philox (uint2 st,
k += 0x9e3779b9u;
}

- return convert_float2(st) / 2147483648.0 - 1.0; + return convert_float2(st) / 2147483648.0f - 1.0f; }

__kernel void kernel_noise (__global float *out, @@ -47,19 +47,19 @@ __kernel void kernel_noise (__global float *out, /* Skew the input point and find the lowest corner of the containing simplex. */

- s = (p.x + p.y) * (sqrt(3.0) - 1) / 2; + s = (p.x + p.y) * (sqrt(3.0f) - 1) / 2; i = floor(p + s);

/* Calculate the (unskewed) distance between the input point and all simplex corners. */

- s = (i.x + i.y) * (3 - sqrt(3.0)) / 6; + s = (i.x + i.y) * (3 - sqrt(3.0f)) / 6; u[0] = p - i + s;

di = u[0].x >= u[0].y ? (float2)(1, 0) : (float2)(0, 1);
- u[1] = u[0] - di + (3 - sqrt(3.0)) / 6; - u[2] = u[0] - 1 + (3 - sqrt(3.0)) / 3; + u[1] = u[0] - di + (3 - sqrt(3.0f)) / 6; + u[2] = u[0] - 1 + (3 - sqrt(3.0f)) / 3;
/* Calculate gradients for each corner vertex. We convert to * signed int first to avoid implementation-defined behavior for @@ -72,7 +72,7 @@ __kernel void kernel_noise (__global float *out,
for (k = 0, n = 0 ; k < 3 ; k += 1) {
- t = 0.5 - dot(u[k], u[k]); + t = 0.5f - dot(u[k], u[k]);
if (t > 0)
{

1.9.0
Jan Vesely
2014-06-05 13:09:41 UTC (almost 10 years ago)

[PATCH 1/1] cl: Make noise-simplex kernel OpenCL 1.1 complaint

gentle ping

On Sun, 2014-05-18 at 16:38 -0400, Jan Vesely wrote:

No static keyword, f suffix for float constants Fixes one error and 7 warnings

Signed-off-by: Jan Vesely ---
Makes the op run on mesa/clover.

opencl/noise-simplex.cl | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-)

diff --git a/opencl/noise-simplex.cl b/opencl/noise-simplex.cl index bb59758..ce8aa40 100644
--- a/opencl/noise-simplex.cl
+++ b/opencl/noise-simplex.cl
@@ -1,6 +1,6 @@
#define MAX_RANK 3

-static float2
+float2
philox (uint2 st,
uint k)
{
@@ -17,7 +17,7 @@ philox (uint2 st,
k += 0x9e3779b9u;
}

- return convert_float2(st) / 2147483648.0 - 1.0; + return convert_float2(st) / 2147483648.0f - 1.0f; }

__kernel void kernel_noise (__global float *out, @@ -47,19 +47,19 @@ __kernel void kernel_noise (__global float *out, /* Skew the input point and find the lowest corner of the containing simplex. */

- s = (p.x + p.y) * (sqrt(3.0) - 1) / 2; + s = (p.x + p.y) * (sqrt(3.0f) - 1) / 2; i = floor(p + s);

/* Calculate the (unskewed) distance between the input point and all simplex corners. */

- s = (i.x + i.y) * (3 - sqrt(3.0)) / 6; + s = (i.x + i.y) * (3 - sqrt(3.0f)) / 6; u[0] = p - i + s;

di = u[0].x >= u[0].y ? (float2)(1, 0) : (float2)(0, 1);
- u[1] = u[0] - di + (3 - sqrt(3.0)) / 6; - u[2] = u[0] - 1 + (3 - sqrt(3.0)) / 3; + u[1] = u[0] - di + (3 - sqrt(3.0f)) / 6; + u[2] = u[0] - 1 + (3 - sqrt(3.0f)) / 3;
/* Calculate gradients for each corner vertex. We convert to * signed int first to avoid implementation-defined behavior for @@ -72,7 +72,7 @@ __kernel void kernel_noise (__global float *out,
for (k = 0, n = 0 ; k < 3 ; k += 1) {
- t = 0.5 - dot(u[k], u[k]); + t = 0.5f - dot(u[k], u[k]);
if (t > 0)
{

Jan Vesely 
scl
2014-06-07 20:31:08 UTC (almost 10 years ago)

[PATCH 1/1] cl: Make noise-simplex kernel OpenCL 1.1 complaint

On 5.6.2014 at 3:09 PM Jan Vesely wrote:

gentle ping

On Sun, 2014-05-18 at 16:38 -0400, Jan Vesely wrote:

No static keyword, f suffix for float constants Fixes one error and 7 warnings

Hi Jan,

first of all thank you for your work! I'm more active in the GIMP project than in GEGL, but I appreciate your efforts and would like to push it a bit further, so here we go:

- Is your code tested or can you provide test drivers that fit into the GEGL codebase? - I could try to test the code, but I only have Linux in a Virtualbox VM. Can you tell whether Mesa/Gallium3d/OpenCL would work in such an environment?
- Are there any side effects of the patch known (especially from making the static variable a non-static one)?

Kind regards,

Sven

Jan Vesely
2014-06-07 21:19:59 UTC (almost 10 years ago)

[PATCH 1/1] cl: Make noise-simplex kernel OpenCL 1.1 complaint

Hi,

thank you for your response.

On Sat, 2014-06-07 at 22:31 +0200, scl wrote:

On 5.6.2014 at 3:09 PM Jan Vesely wrote:

gentle ping

On Sun, 2014-05-18 at 16:38 -0400, Jan Vesely wrote:

No static keyword, f suffix for float constants Fixes one error and 7 warnings

Hi Jan,

first of all thank you for your work! I'm more active in the GIMP project than in GEGL, but I appreciate your efforts and would like to push it a bit further, so here we go:

- Is your code tested or can you provide test drivers that fit into the GEGL codebase?

yes, the noise-simplex.xml test passes on my AMD TURKS gpu (Radeon HD 7570).
I'm not sure how to provide a test driver. I also ran the gegl testsuite on Intel OCL SDK and this patch does not change any result (all tests pass)

- I could try to test the code, but I only have Linux in a Virtualbox VM. Can you tell whether Mesa/Gallium3d/OpenCL would work in such an environment?

It won't. VBox does not use mesa at all. you'd need linux running on a machine with amd gpu to test it, as well as recent versions of mesa, clang, llvm, and libclc .
It should be possible to test ocl 1.1 syntax using clang CLC frontend but I have not investigated it.

- Are there any side effects of the patch known (especially from making the static variable a non-static one)?

No side effects.
static keyword is removed from a function not variable (although it's not supported for either in OCL 1.1)

this patch is in line with previous modifications to make ocl kernels ocl 1.1 complaint:
https://git.gnome.org/browse/gegl/commit/?id=75fc532a561ddcf3f4d50c26b94aa31b047abfa9 https://git.gnome.org/browse/gegl/commit/?id=1fb0a80f236d8ed75146e57e568b19ff1d906d99 https://git.gnome.org/browse/gegl/commit/?id=bd9f98e2b0c4add1325a62a88260fbd01cab6143

the only difference is that this time the associated .h file is not in the repository (not sure if it's intentional)

thank you, Jan

Kind regards,

Sven

_______________________________________________ gegl-developer-list mailing list
List address: gegl-developer-list@gnome.org List membership: https://mail.gnome.org/mailman/listinfo/gegl-developer-list

Jan Vesely 
scl
2014-06-14 09:52:33 UTC (almost 10 years ago)

[PATCH 1/1] cl: Make noise-simplex kernel OpenCL 1.1 complaint

Hi,

I've built it successfully on Linux Debian Testing (32 and 64 bit). For more information see also the GEGL branch wip/noise-simplex-cl-patch and https://build.gimp.org/job/gegl-noise-simplex-cl/

@Pippin, Daniel, Victor: can you please have a look on this patch? If you have no objections, then I would merge it to GEGL master in a week and remove the temporary GEGL branch and Jenkins job then.

Kind regards,

Sven

Daniel Sabo
2014-06-18 07:22:20 UTC (almost 10 years ago)

[PATCH 1/1] cl: Make noise-simplex kernel OpenCL 1.1 complaint

This all looks "obviously correct", we've cleaned up similar things before because most OpenCL implementations will fudge around these things but some complain. I have not been involved enough with GEGL for a while to go pushing stuff; you should add massimo to your list of people to poke on IRC about this, he has a good eye for reviewing these and has the AMD OpenCL implementation set up to test on.

- Daniel

On Sat, Jun 14, 2014 at 2:52 AM, scl wrote:

Hi,

I've built it successfully on Linux Debian Testing (32 and 64 bit). For more information see also the GEGL branch wip/noise-simplex-cl-patch and https://build.gimp.org/job/gegl-noise-simplex-cl/

@Pippin, Daniel, Victor: can you please have a look on this patch? If you have no objections, then I would merge it to GEGL master in a week and remove the temporary GEGL branch and Jenkins job then.

Kind regards,

Sven

scl
2014-06-19 19:08:40 UTC (almost 10 years ago)

[PATCH 1/1] cl: Make noise-simplex kernel OpenCL 1.1 complaint

Hi Jan,

my local 'make check' tests passed and the GEGL developers had no objections, so I've committed your patch to GEGL master.

Thank you for your contribution and feel free to continue your work!

Greetings,

Sven