[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.
[PATCH 1/1] cl: Make noise-simplex kernel OpenCL 1.1 complaint | Jan Vesely | 18 May 20:38 |
[PATCH 1/1] cl: Make noise-simplex kernel OpenCL 1.1 complaint | Jan Vesely | 05 Jun 13:09 |
[PATCH 1/1] cl: Make noise-simplex kernel OpenCL 1.1 complaint | scl | 07 Jun 20:31 |
[PATCH 1/1] cl: Make noise-simplex kernel OpenCL 1.1 complaint | Jan Vesely | 07 Jun 21:19 |
[PATCH 1/1] cl: Make noise-simplex kernel OpenCL 1.1 complaint | scl | 14 Jun 09:52 |
[PATCH 1/1] cl: Make noise-simplex kernel OpenCL 1.1 complaint | Daniel Sabo | 18 Jun 07:22 |
[PATCH 1/1] cl: Make noise-simplex kernel OpenCL 1.1 complaint | scl | 19 Jun 19:08 |
[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
[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
[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
[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
[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
[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
[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