On Thu, 2014-06-19 at 17:27 +0200, Bruno Jimenez wrote: > Hi, > > Thanks for catching this bug! > > Reviewed-by: Bruno Jiménez <brunoji...@gmail.com> > > Also, could you please send me a copy of the tests?
This got triggered by one of the gegl test. You'll need the attached
patches to run gegl git[0] on mesa/clover/r600. you 'll need babl git[1]
too. Let me know if I can help with setting it up.
Sry, I don't have any smaller reproducer.
Jan
[0] git://git.gnome.org/gegl
[1] git://git.gnome.org/babl
>
> On Thu, 2014-06-19 at 10:21 -0400, Jan Vesely wrote:
> > The important part is the change of the condition to <= 0. Otherwise the
> > loop
> > gets stuck never actually growing the pool.
> >
> > The change in the aux-need calculation guarantees max 2 iterations, and
> > avoids wasting memory in case a smaller item can't fit into a relatively
> > larger
> > pool.
> >
> > Signed-off-by: Jan Vesely <jan.ves...@rutgers.edu>
> > CC: Bruno Jimenez <brunoji...@gmail.com>
> > ---
> >
> > This fixes hang in gegl colors.xml test
> >
> > src/gallium/drivers/r600/compute_memory_pool.c | 7 +++++--
> > 1 file changed, 5 insertions(+), 2 deletions(-)
> >
> > diff --git a/src/gallium/drivers/r600/compute_memory_pool.c
> > b/src/gallium/drivers/r600/compute_memory_pool.c
> > index ec8c470..0b6d2da6 100644
> > --- a/src/gallium/drivers/r600/compute_memory_pool.c
> > +++ b/src/gallium/drivers/r600/compute_memory_pool.c
> > @@ -320,8 +320,11 @@ int compute_memory_finalize_pending(struct
> > compute_memory_pool* pool,
> > int64_t need = item->size_in_dw+2048 -
> > (pool->size_in_dw - allocated);
> >
> > - if (need < 0) {
> > - need = pool->size_in_dw / 10;
> > + if (need <= 0) {
> > + /* There's enough free space, but it's too
> > + * fragmented. Assume half of the item can fit
> > + * int the last chunk */
> > + need = (item->size_in_dw / 2) + ITEM_ALIGNMENT;
> > }
> >
> > need = align(need, ITEM_ALIGNMENT);
>
>
--
Jan Vesely <jan.ves...@rutgers.edu>
From 762d9251f336e7c0c1db7b1af9c01880477ac6ef Mon Sep 17 00:00:00 2001 From: Jan Vesely <jan.ves...@rutgers.edu> Date: Wed, 13 Nov 2013 11:34:32 -0500 Subject: [PATCH 1/3] opencl: Don't load some unused extensions. These are not available on mesa. Signed-off-by: Jan Vesely <jan.ves...@rutgers.edu> --- gegl/opencl/gegl-cl-init.c | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/gegl/opencl/gegl-cl-init.c b/gegl/opencl/gegl-cl-init.c index 8ae1ef1..5024421 100644 --- a/gegl/opencl/gegl-cl-init.c +++ b/gegl/opencl/gegl-cl-init.c @@ -501,9 +501,9 @@ gegl_cl_init_load_functions (GError **error) CL_LOAD_FUNCTION (clReleaseContext) CL_LOAD_FUNCTION (clReleaseMemObject) - CL_LOAD_FUNCTION (clCreateFromGLTexture2D) - CL_LOAD_FUNCTION (clEnqueueAcquireGLObjects) - CL_LOAD_FUNCTION (clEnqueueReleaseGLObjects) +// CL_LOAD_FUNCTION (clCreateFromGLTexture2D) +// CL_LOAD_FUNCTION (clEnqueueAcquireGLObjects) +// CL_LOAD_FUNCTION (clEnqueueReleaseGLObjects) return TRUE; } -- 1.9.3
From 821c1bbc2bccce3b6706769546fa0d5ddc7d83d6 Mon Sep 17 00:00:00 2001
From: Jan Vesely <jan.ves...@rutgers.edu>
Date: Wed, 18 Dec 2013 18:19:16 -0500
Subject: [PATCH 2/3] Don't use color_kernels, they are broken on mesa
Signed-off-by: Jan Vesely <jan.ves...@rutgers.edu>
---
gegl/opencl/gegl-cl-color.c | 7 +++++--
1 file changed, 5 insertions(+), 2 deletions(-)
diff --git a/gegl/opencl/gegl-cl-color.c b/gegl/opencl/gegl-cl-color.c
index fbcd8ab..ac693f6 100644
--- a/gegl/opencl/gegl-cl-color.c
+++ b/gegl/opencl/gegl-cl-color.c
@@ -232,8 +232,11 @@ gegl_cl_color_supported (const Babl *in_format,
if (in_format == out_format)
return GEGL_CL_COLOR_EQUAL;
- if (color_kernels_hash && find_color_kernel (in_format, out_format))
- return GEGL_CL_COLOR_CONVERT;
+ if (color_kernels_hash && find_color_kernel (in_format, out_format)) {
+// GEGL_NOTE (GEGL_DEBUG_OPENCL, "Found OpenCL conversion from: %s -> %s",
+// babl_get_name(in_format), babl_get_name(out_format));
+// return GEGL_CL_COLOR_CONVERT;
+ }
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Missing OpenCL conversion for %s -> %s",
babl_get_name (in_format),
--
1.9.3
From dc909e2ba904aaa7bde34a49e59a36f93f1ac31b Mon Sep 17 00:00:00 2001
From: Jan Vesely <jan.ves...@rutgers.edu>
Date: Fri, 16 May 2014 19:39:06 -0400
Subject: [PATCH 3/3] cl: Make noise-simplex kernel OpenCL 1.1 complaint
Fixes one error and 7 warnings
no static keyword, f suffix for float constants
Signed-off-by: Jan Vesely <jan.ves...@rutgers.edu>
---
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.3
signature.asc
Description: This is a digitally signed message part
_______________________________________________ mesa-dev mailing list mesa-dev@lists.freedesktop.org http://lists.freedesktop.org/mailman/listinfo/mesa-dev
