[gegl] opencl: color kernels in new format and cosmetic changes
- From: Victor Matheus de Araujo Oliveira <vmaolive src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] opencl: color kernels in new format and cosmetic changes
- Date: Sun, 13 Jan 2013 01:27:58 +0000 (UTC)
commit 96a0494192261ab5f1a0fe1297526c0a10c2cfb4
Author: Victor Oliveira <victormatheus gmail com>
Date: Sat Jan 12 23:27:40 2013 -0200
opencl: color kernels in new format and cosmetic changes
gegl/buffer/gegl-buffer-cl-cache.c | 6 +-
gegl/buffer/gegl-buffer-cl-iterator.c | 55 +--
gegl/opencl/Makefile.am | 1 -
gegl/opencl/gegl-cl-color-kernel.h | 637 ------------------------
gegl/opencl/gegl-cl-color.c | 49 +-
gegl/opencl/gegl-cl-init.c | 1 +
gegl/opencl/gegl-cl.h | 4 +
gegl/operation/gegl-operation-point-composer.c | 19 +-
gegl/operation/gegl-operation-point-filter.c | 16 +-
operations/common/Makefile.am | 2 +-
10 files changed, 79 insertions(+), 711 deletions(-)
---
diff --git a/gegl/buffer/gegl-buffer-cl-cache.c b/gegl/buffer/gegl-buffer-cl-cache.c
index 8c7fca0..621c054 100644
--- a/gegl/buffer/gegl-buffer-cl-cache.c
+++ b/gegl/buffer/gegl-buffer-cl-cache.c
@@ -126,8 +126,6 @@ gegl_buffer_cl_cache_new (GeglBuffer *buffer,
g_static_mutex_unlock (&cache_mutex);
}
-#define CL_ERROR {GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in %s:%d %s - %s\n", __FILE__, __LINE__, __func__, gegl_cl_errstring(cl_err)); goto error;}
-
gboolean
gegl_buffer_cl_cache_flush2 (GeglTileHandlerCache *cache,
const GeglRectangle *roi)
@@ -171,7 +169,7 @@ gegl_buffer_cl_cache_flush2 (GeglTileHandlerCache *cache,
if (need_cl)
{
cl_err = gegl_clFinish (gegl_cl_get_command_queue ());
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
g_static_mutex_lock (&cache_mutex);
@@ -262,5 +260,3 @@ gegl_buffer_cl_cache_invalidate (GeglBuffer *buffer,
#endif
}
-
-#undef CL_ERROR
diff --git a/gegl/buffer/gegl-buffer-cl-iterator.c b/gegl/buffer/gegl-buffer-cl-iterator.c
index 8f36975..4daed6d 100644
--- a/gegl/buffer/gegl-buffer-cl-iterator.c
+++ b/gegl/buffer/gegl-buffer-cl-iterator.c
@@ -36,8 +36,6 @@
#include "opencl/gegl-cl.h"
-#define CL_ERROR {GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in %s:%d %s - %s\n", __FILE__, __LINE__, __func__, gegl_cl_errstring(cl_err)); goto error;}
-
typedef struct GeglBufferClIterators
{
/* current region of interest */
@@ -206,6 +204,7 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
gboolean result = FALSE;
gint no, j;
cl_int cl_err = 0;
+ int color_err = 0;
if (i->is_finished)
g_error ("%s called on finished buffer iterator", G_STRFUNC);
@@ -246,9 +245,9 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
if (i->conv[no] == GEGL_CL_COLOR_CONVERT)
for (j=0; j < i->n; j++)
{
- cl_err = gegl_cl_color_conv (i->tex_op[no][j], i->tex_buf[no][j], i->size[no][j],
- i->format[no], i->buffer[no]->soft_format);
- if (cl_err == FALSE) CL_ERROR;
+ color_err = gegl_cl_color_conv (i->tex_op[no][j], i->tex_buf[no][j], i->size[no][j],
+ i->format[no], i->buffer[no]->soft_format);
+ if (color_err) goto error;
}
/* GPU -> CPU */
@@ -265,7 +264,7 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
i->tex_op[no][j], CL_TRUE,
0, i->size[no][j] * i->op_cl_format_size [no], data,
0, NULL, NULL);
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
/* color conversion using BABL */
gegl_buffer_set (i->buffer[no], &i->roi[no][j], 0, i->format[no], data, GEGL_AUTO_ROWSTRIDE);
@@ -285,14 +284,14 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
CL_MAP_READ,
0, i->size[no][j] * i->buf_cl_format_size [no],
0, NULL, NULL, &cl_err);
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
/* color conversion using BABL */
gegl_buffer_set (i->buffer[no], &i->roi[no][j], i->format[no], data, GEGL_AUTO_ROWSTRIDE);
cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_buf[no][j], data,
0, NULL, NULL);
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
}
#endif
}
@@ -301,7 +300,7 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
/* Run! */
cl_err = gegl_clFinish(gegl_cl_get_command_queue());
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
for (no=0; no < i->iterators; no++)
for (j=0; j < i->n; j++)
@@ -361,21 +360,21 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY,
i->size[no][j] * i->op_cl_format_size [no],
NULL, &cl_err);
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
/* pre-pinned memory */
data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_op[no][j], CL_TRUE,
CL_MAP_WRITE,
0, i->size[no][j] * i->op_cl_format_size [no],
0, NULL, NULL, &cl_err);
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
/* color conversion using BABL */
gegl_buffer_get (i->buffer[no], &i->roi[no][j], 1.0, i->format[no], data, GEGL_AUTO_ROWSTRIDE, GEGL_ABYSS_NONE);
cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_op[no][j], data,
0, NULL, NULL);
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
i->tex[no][j] = i->tex_op[no][j];
@@ -398,21 +397,21 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY,
i->size[no][j] * i->buf_cl_format_size [no],
NULL, &cl_err);
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
/* pre-pinned memory */
data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_buf[no][j], CL_TRUE,
CL_MAP_WRITE,
0, i->size[no][j] * i->buf_cl_format_size [no],
0, NULL, NULL, &cl_err);
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
/* color conversion will be performed in the GPU later */
gegl_buffer_get (i->buffer[no], &i->roi[no][j], 1.0, i->buffer[no]->soft_format, data, GEGL_AUTO_ROWSTRIDE, GEGL_ABYSS_NONE);
cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_buf[no][j], data,
0, NULL, NULL);
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
}
i->tex[no][j] = i->tex_buf[no][j];
@@ -436,21 +435,21 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY,
i->size[no][j] * i->buf_cl_format_size [no],
NULL, &cl_err);
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
/* pre-pinned memory */
data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_buf[no][j], CL_TRUE,
CL_MAP_WRITE,
0, i->size[no][j] * i->buf_cl_format_size [no],
0, NULL, NULL, &cl_err);
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
/* color conversion will be performed in the GPU later */
gegl_buffer_get (i->buffer[no], &i->roi[no][j], 1.0, i->buffer[no]->soft_format, data, GEGL_AUTO_ROWSTRIDE, GEGL_ABYSS_NONE);
cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_buf[no][j], data,
0, NULL, NULL);
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
}
g_assert (i->tex_op[no][j] == NULL);
@@ -458,13 +457,13 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
CL_MEM_READ_WRITE,
i->size[no][j] * i->op_cl_format_size [no],
NULL, &cl_err);
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
/* color conversion in the GPU (input) */
g_assert (i->tex_buf[no][j] && i->tex_op[no][j]);
- cl_err = gegl_cl_color_conv (i->tex_buf[no][j], i->tex_op[no][j], i->size[no][j],
- i->buffer[no]->soft_format, i->format[no]);
- if (cl_err == FALSE) CL_ERROR;
+ color_err = gegl_cl_color_conv (i->tex_buf[no][j], i->tex_op[no][j], i->size[no][j],
+ i->buffer[no]->soft_format, i->format[no]);
+ if (color_err) goto error;
i->tex[no][j] = i->tex_op[no][j];
@@ -487,7 +486,7 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
CL_MEM_WRITE_ONLY,
i->size[no][j] * i->op_cl_format_size [no],
NULL, &cl_err);
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
i->tex[no][j] = i->tex_op[no][j];
@@ -502,7 +501,7 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
CL_MEM_READ_WRITE, /* cache */
i->size[no][j] * i->buf_cl_format_size [no],
NULL, &cl_err);
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
i->tex[no][j] = i->tex_buf[no][j];
@@ -517,14 +516,14 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
CL_MEM_READ_WRITE, /* cache */
i->size[no][j] * i->buf_cl_format_size [no],
NULL, &cl_err);
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
g_assert (i->tex_op[no][j] == NULL);
i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_READ_WRITE,
i->size[no][j] * i->op_cl_format_size [no],
NULL, &cl_err);
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
i->tex[no][j] = i->tex_op[no][j];
@@ -542,7 +541,7 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
CL_MEM_READ_WRITE,
i->size[no][j] * i->op_cl_format_size [no],
NULL, &cl_err);
- if (cl_err != CL_SUCCESS) CL_ERROR;
+ CL_CHECK;
i->tex[no][j] = i->tex_op[no][j];
}
@@ -616,5 +615,3 @@ gegl_buffer_cl_iterator_new (GeglBuffer *buffer,
gegl_buffer_cl_iterator_add (i, buffer, roi, format, flags, GEGL_ABYSS_NONE);
return i;
}
-
-#undef CL_ERROR
diff --git a/gegl/opencl/Makefile.am b/gegl/opencl/Makefile.am
index 420df54..87a896a 100644
--- a/gegl/opencl/Makefile.am
+++ b/gegl/opencl/Makefile.am
@@ -32,7 +32,6 @@ libcl_sources = \
gegl-cl-types.h \
gegl-cl-init.c \
gegl-cl-init.h \
- gegl-cl-color-kernel.h \
gegl-cl-color.c \
gegl-cl-color.h
diff --git a/gegl/opencl/gegl-cl-color.c b/gegl/opencl/gegl-cl-color.c
index 6a202db..aed1d77 100644
--- a/gegl/opencl/gegl-cl-color.c
+++ b/gegl/opencl/gegl-cl-color.c
@@ -26,10 +26,10 @@
#include "gegl.h"
#include "gegl/gegl-debug.h"
+#include "gegl-cl.h"
#include "gegl-cl-color.h"
-#include "gegl-cl-init.h"
-#include "gegl-cl-color-kernel.h"
+#include "opencl/colors.cl.h"
static GeglClRunData *kernels_color = NULL;
@@ -145,7 +145,7 @@ gegl_cl_color_compile_kernels(void)
format[9] = babl_format ("R'G'B'A u8");
format[10] = babl_format ("R'G'B' u8");
- kernels_color = gegl_cl_compile_and_build (kernel_color_source, kernel_name);
+ kernels_color = gegl_cl_compile_and_build (colors_cl_source, kernel_name);
}
@@ -275,8 +275,6 @@ gegl_cl_color_supported (const Babl *in_format,
return GEGL_CL_COLOR_NOT_SUPPORTED;
}
-#define CL_ERROR {GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in %s:%d %s - %s\n", __FILE__, __LINE__, __func__, gegl_cl_errstring(errcode)); return FALSE;}
-
gboolean
gegl_cl_color_conv (cl_mem in_tex,
cl_mem out_tex,
@@ -284,7 +282,7 @@ gegl_cl_color_conv (cl_mem in_tex,
const Babl *in_format,
const Babl *out_format)
{
- int errcode;
+ cl_int cl_err;
if (gegl_cl_color_supported (in_format, out_format) == GEGL_CL_COLOR_NOT_SUPPORTED)
return FALSE;
@@ -295,35 +293,36 @@ gegl_cl_color_conv (cl_mem in_tex,
gegl_cl_color_babl (in_format, &s);
/* just copy in_tex to out_tex */
- errcode = gegl_clEnqueueCopyBuffer (gegl_cl_get_command_queue(),
- in_tex, out_tex, 0, 0, size * s,
- 0, NULL, NULL);
- if (errcode != CL_SUCCESS) CL_ERROR
+ cl_err = gegl_clEnqueueCopyBuffer (gegl_cl_get_command_queue(),
+ in_tex, out_tex, 0, 0, size * s,
+ 0, NULL, NULL);
+ CL_CHECK;
- errcode = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
- if (errcode != CL_SUCCESS) CL_ERROR
+ cl_err = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
+ CL_CHECK;
}
else
{
gint k = choose_kernel (in_format, out_format);
- errcode = gegl_clSetKernelArg(kernels_color->kernel[k], 0, sizeof(cl_mem), (void*)&in_tex);
- if (errcode != CL_SUCCESS) CL_ERROR
+ cl_err = gegl_clSetKernelArg(kernels_color->kernel[k], 0, sizeof(cl_mem), (void*)&in_tex);
+ CL_CHECK;
- errcode = gegl_clSetKernelArg(kernels_color->kernel[k], 1, sizeof(cl_mem), (void*)&out_tex);
- if (errcode != CL_SUCCESS) CL_ERROR
+ cl_err = gegl_clSetKernelArg(kernels_color->kernel[k], 1, sizeof(cl_mem), (void*)&out_tex);
+ CL_CHECK;
- errcode = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
- kernels_color->kernel[k], 1,
- NULL, &size, NULL,
- 0, NULL, NULL);
- if (errcode != CL_SUCCESS) CL_ERROR
+ cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+ kernels_color->kernel[k], 1,
+ NULL, &size, NULL,
+ 0, NULL, NULL);
+ CL_CHECK;
- errcode = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
- if (errcode != CL_SUCCESS) CL_ERROR
+ cl_err = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
+ CL_CHECK;
}
+ return FALSE;
+
+error:
return TRUE;
}
-
-#undef CL_ERROR
diff --git a/gegl/opencl/gegl-cl-init.c b/gegl/opencl/gegl-cl-init.c
index c8dafbf..f97c9cc 100644
--- a/gegl/opencl/gegl-cl-init.c
+++ b/gegl/opencl/gegl-cl-init.c
@@ -302,6 +302,7 @@ gegl_cl_init (GError **error)
err = gegl_clGetDeviceIDs (cl_state.platform, CL_DEVICE_TYPE_DEFAULT, 1, &cl_state.device, NULL);
if(err != CL_SUCCESS)
{
+ GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", gegl_cl_errstring(err));
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Could not create device");
return FALSE;
}
diff --git a/gegl/opencl/gegl-cl.h b/gegl/opencl/gegl-cl.h
index f7a2641..b70fd27 100644
--- a/gegl/opencl/gegl-cl.h
+++ b/gegl/opencl/gegl-cl.h
@@ -23,4 +23,8 @@
#include "gegl-cl-init.h"
#include "gegl-cl-color.h"
+#define CL_ERROR {GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in %s:%d %s - %s\n", __FILE__, __LINE__, __func__, gegl_cl_errstring(cl_err)); goto error;}
+
+#define CL_CHECK {if (cl_err != CL_SUCCESS) CL_ERROR;}
+
#endif
diff --git a/gegl/operation/gegl-operation-point-composer.c b/gegl/operation/gegl-operation-point-composer.c
index 5401e76..02fb0e3 100644
--- a/gegl/operation/gegl-operation-point-composer.c
+++ b/gegl/operation/gegl-operation-point-composer.c
@@ -189,6 +189,7 @@ gegl_operation_point_composer_cl_process (GeglOperation *operation,
cl_err = point_composer_class->cl_process(operation, i->tex[read][j],
(aux)? i->tex[foo][j] : NULL,
i->tex[0][j], i->size[0][j], &i->roi[0][j], level);
+ CL_CHECK;
}
else if (operation_class->cl_data)
{
@@ -196,36 +197,40 @@ gegl_operation_point_composer_cl_process (GeglOperation *operation,
GeglClRunData *cl_data = operation_class->cl_data;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[read][j]);
+ CL_CHECK;
if (aux)
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[foo][j]);
else
- cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, 0, 0);
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), NULL);
+ CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[0][j]);
+ CL_CHECK;
gegl_operation_cl_set_kernel_args (operation, cl_data->kernel[0], &p, &cl_err);
+ CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 1,
NULL, &i->size[0][j], NULL,
0, NULL, NULL);
+ CL_CHECK;
}
else
{
g_warning ("OpenCL support enabled, but no way to execute");
return FALSE;
}
-
- if (cl_err != CL_SUCCESS)
- {
- GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in GeglOperationPointComposer Kernel: %s", gegl_cl_errstring(cl_err));
- return FALSE;
- }
}
}
}
+
return TRUE;
+
+error:
+ GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in GeglOperationPointComposer Kernel: %s", gegl_cl_errstring(cl_err));
+ return FALSE;
}
static gboolean
diff --git a/gegl/operation/gegl-operation-point-filter.c b/gegl/operation/gegl-operation-point-filter.c
index 7d55f2e..327feb4 100644
--- a/gegl/operation/gegl-operation-point-filter.c
+++ b/gegl/operation/gegl-operation-point-filter.c
@@ -115,6 +115,7 @@ gegl_operation_point_filter_cl_process (GeglOperation *operation,
{
cl_err = point_filter_class->cl_process(operation, i->tex[read][j], i->tex[0][j],
i->size[0][j], &i->roi[0][j], level);
+ CL_CHECK;
}
else if (operation_class->cl_data)
{
@@ -122,30 +123,33 @@ gegl_operation_point_filter_cl_process (GeglOperation *operation,
GeglClRunData *cl_data = operation_class->cl_data;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[read][j]);
+ CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[ 0 ][j]);
+ CL_CHECK;
gegl_operation_cl_set_kernel_args (operation, cl_data->kernel[0], &p, &cl_err);
+ CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 1,
NULL, &i->size[0][j], NULL,
0, NULL, NULL);
+ CL_CHECK;
}
else
{
g_warning ("OpenCL support enabled, but no way to execute");
return FALSE;
}
-
- if (cl_err != CL_SUCCESS)
- {
- GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in GeglOperationPointFilter Kernel: %s", gegl_cl_errstring(cl_err));
- return FALSE;
- }
}
}
}
+
return TRUE;
+
+error:
+ GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in GeglOperationPointComposer Kernel: %s", gegl_cl_errstring(cl_err));
+ return FALSE;
}
static gboolean
diff --git a/operations/common/Makefile.am b/operations/common/Makefile.am
index c106c3b..622c103 100644
--- a/operations/common/Makefile.am
+++ b/operations/common/Makefile.am
@@ -2,5 +2,5 @@ SUBDIRS = perlin
include $(top_srcdir)/operations/Makefile-operations.am
AM_CPPFLAGS += \
+ -I$(top_srcdir) \
-I$(top_srcdir)/operations/common
- -I$(top_srcdir)/opencl
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
