Skip to content

Commit

Permalink
Merge pull request #7 from clEsperanto/development
Browse files Browse the repository at this point in the history 
Merge to prepare CLIJ2 release
  • Loading branch information
@haesleinhuepf
haesleinhuepf authored Jun 12, 2020
2 parents 8bb15df + 2fc5ff2 commit 9546d0f
Show file tree
Hide file tree
Showing 337 changed files with 9,308 additions and 4,485 deletions.
128 changes: 38 additions & 90 deletions README.md
View file
Original file line number Diff line number Diff line change
@@ -1,7 +1,7 @@
# clij-opencl-kernels

This repository contains a collection of [OpenCL](https://www.khronos.org/opencl/) [kernels for generic image
processing](https://github.com/clij/clij-opencl-kernels/tree/master/src/main/java/net/haesleinhuepf/clij/kernels).
This repository contains a collection of [OpenCL](https://www.khronos.org/opencl/) kernels for [image
processing](https://github.com/clij/clij-opencl-kernels/tree/development/src/main/java/net/haesleinhuepf/clij/kernels).
The [CLIJ](https://clij.github.io) is build on top of it allowing
[ImageJ](https://imagej.nih.gov/ij/) / [Fiji](https://fiji.sc) users in doing
GPU-accelerated image processing without the need for learning OpenCL.
Expand All @@ -11,7 +11,7 @@ GPU-accelerated image processing without the need for learning OpenCL.
Robert Haase, Loic Alain Royer, Peter Steinbach, Deborah Schmidt,
Alexandr Dibrov, Uwe Schmidt, Martin Weigert, Nicola Maghelli, Pavel Tomancak,
Florian Jug, Eugene W Myers.
*CLIJ: GPU-accelerated image processing for everyone*. BioRxiv preprint. [https://doi.org/10.1101/660704](https://doi.org/10.1101/660704)
*CLIJ: GPU-accelerated image processing for everyone*. [Nat Methods (2019) doi:10.1038/s41592-019-0650-1](https://doi.org/10.1038/s41592-019-0650-1)

## Why a custom OpenCL-dialect?

Expand All @@ -20,10 +20,10 @@ Theoretically, one has to write OpenCL-kernels specifically for given input- and
adding images of type `float` resulting in a `float` image and a kernel for adding image of type `uint8` resulting in
an image of type `float`. Furthermore, OpenCL defines images and buffers. However, as both are arrays of pixel
intensities in memory, we wanted to access them in a unified way. As this would result in a ridiculous large number of individual kernel implementations, we used
a dialect where placeholders such as `DTYPE_OUT` represent the pixel type of the output image.
a dialect where placeholders such as `IMAGE_src_PIXEL_TYPE` represent the pixel type of the output image.

## List of placeholders
The following list of placeholders are used at the moment:
The following list of placeholders are used at the moment. The name `imagename` must contain `src` or `dst` to differentiate `readonly` and `writeonly` images.
<table border="1">

<tr>
Expand All @@ -32,8 +32,8 @@ The following list of placeholders are used at the moment:
</tr>

<tr>
<td><pre>CONVERT_DTYPE_IN</pre></td>
<td rowspan="2"><pre>
<td><pre>CONVERT_imagename_PIXEL_TYPE</pre></td>
<td><pre>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L303">clij_convert_char_sat</a>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L292">clij_convert_uchar_sat</a>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L325">clij_convert_short_sat</a>
Expand All @@ -43,106 +43,72 @@ The following list of placeholders are used at the moment:
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L355">clij_convert_float_sat</a>
</pre>
</td>
<td rowspan="2">Convert any number to a given type.</td>
</tr>
<tr>
<td><pre>CONVERT_DTYPE_OUT</pre></td>
</tr>
<tr>
<td><pre>DTYPE_IMAGE_IN_2D</pre></td>
<td><pre>
__read_only image2d_t
__global char*
__global uchar*
__global short*
__global ushort*
__global float*
</pre></td>
<td>Two dimensional input image type definition</td>
<td>Convert any number to a given type.</td>
</tr>
<tr>
<td><pre>DTYPE_IMAGE_IN_3D</pre></td>
<td><pre>IMAGE_imagename_TYPE</pre></td>
<td><pre>
__read_only image3d_t
__global char*
__global uchar*
__global short*
__global ushort*
__global float*
</pre></td>
<td>Three dimensional input image type definition</td>
</tr>
<tr>
<td><pre>DTYPE_IMAGE_OUT_2D</pre></td>
<td><pre>
__read_only image2d_t
__write_only image2d_t
__global char*
__global uchar*
__global short*
__global ushort*
__global float*
</pre></td>
<td>Two dimensional output image type definition</td>
</tr>
<tr>
<td><pre>DTYPE_IMAGE_OUT_3D</pre></td>
<td><pre>
__write_only image3d_t
__global char*
__global uchar*
__global short*
__global ushort*
__global float*
</pre></td>
<td>Three dimensional output image type definition</td>
<td>Two dimensional input image type definition. </td>
</tr>
<tr>
<td><pre>DTYPE_IN</pre></td>
<td rowspan="2"><pre>
<td><pre>IMAGE_imagename_PIXEL_TYPE</pre></td>
<td><pre>
char
uchar
short
ushort
float
</pre></td>
<td rowspan="2">Pixel type definition</td>
</tr>
<tr>
<td><pre>DTYPE_OUT</pre></td>
<td>Pixel type definition</td>
</tr>
<tr>
<td><pre>GET_IMAGE_DEPTH</pre></td>
<td><pre>GET_IMAGE_DEPTH(imagename)</pre></td>
<td>constant number</td>
<td>Image size in Z</td>
</tr>
<tr>
<td><pre>GET_IMAGE_HEIGHT</pre></td>
<td><pre>GET_IMAGE_HEIGHT(imagename)</pre></td>
<td>constant number</td>
<td>Image size in Y</td>
</tr>
<tr>
<td><pre>GET_IMAGE_WIDTH</pre></td>
<td><pre>GET_IMAGE_WIDTH(imagename)</pre></td>
<td>constant number</td>
<td>Image size in X</td>
</tr>
<tr>
<td><pre>READ_IMAGE_2D</pre></td>
<td><pre>POS_imagename_TYPE</pre></td>
<td>int2<br>int4</td>
<td>Type of coordinate</td>
</tr>
<tr>
<td><pre>POS_imagename_INSTANCE(pos0,pos1,pos2,pos3)</pre></td>
<td>int2(pos0, pos1)<br>int4(pos0,pos1,pos2,pos3)</td>
<td>instantiate variable of coordinate</td>
</tr>

<tr>
<td><pre>READ_imagename_IMAGE</pre></td>
<td><pre>
<a href="https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/read_imagei2d.html">read_imageui</a>
<a href="https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/read_imagef2d.html">read_imagef</a>
<a href="https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/read_imagei2d.html">read_imageui (2d)</a>
<a href="https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/read_imagef2d.html">read_imagef (2d)</a>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L167">read_buffer2dc</a>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L383">read_buffer2duc</a>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L199">read_buffer2di</a>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L215">read_buffer2dui</a>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L231">read_buffer2df</a>
</pre></td>
<td>Read pixel intensity from a given position</td>
</tr>
<tr>
<td><pre>READ_IMAGE_3D</pre></td>
<td><pre>
<a href="https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/read_imagei3d.html">read_imageui</a>
<a href="https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/read_imagef3d.html">read_imagef</a>
<a href="https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/read_imagei3d.html">read_imageui (3d)</a>
<a href="https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/read_imagef3d.html">read_imagef (3d)</a>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L32">read_buffer3dc</a>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L50">read_buffer3duc</a>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L68">read_buffer3di</a>
Expand All @@ -152,23 +118,17 @@ float
<td>Read pixel intensity from a given position</td>
</tr>
<tr>
<td><pre>WRITE_IMAGE_2D</pre></td>
<td><pre>WRITE_imagename_IMAGE</pre></td>
<td><pre>
<a href="https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/write_image2d.html">write_imageui</a>
<a href="https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/write_image2d.html">write_imagef</a>
<a href="https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/write_image2d.html">write_imageui (2d)</a>
<a href="https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/write_image2d.html">write_imagef (2d)</a>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L247">write_buffer2dc</a>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L256">write_buffer2duc</a>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L265">write_buffer2di</a>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L274">write_buffer2dui</a>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L283">write_buffer2df</a>
</pre></td>
<td>Write pixel intensity to a given position</td>
</tr>
<tr>
<td><pre>WRITE_IMAGE_3D</pre></td>
<td><pre>
<a href="https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/write_image3d.html">write_imageui</a>
<a href="https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/write_image3d.html">write_imagef</a>
<a href="https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/write_image3d.html">write_imageui (3d)</a>
<a href="https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/write_image3d.html">write_imagef (3d)</a>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L122">write_buffer3dc</a>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L131">write_buffer3duc</a>
<a href="https://github.com/clij/clij-clearcl/blob/master/src/main/java/net/haesleinhuepf/clij/clearcl/ocllib/preamble/preamble.cl#L140">write_buffer3di</a>
Expand All @@ -181,17 +141,5 @@ float
</table>

## Known issues
* The described OpenCL dialect has a drawback: Input images all need to be of the same type and output images have to be
of the same type. As this assumption might be violated in practice, we are working on an new version of the dialect to
be release in summer 2020.
* Some of the placeholders represent image size in x/y/z. This results in performance drop in case many images of
different size are processed as individual kernels need to be implemented. We have this in mind for the summer 2020
update of the OpenCL-dialect.
* Image dimensionality is limited to three dimensions.

## Deprecation note
The opencl-kernels presented here will be replaced by Summer 2020. Background
is an upcoming update of the used OpenCL dialect. If you plan to build your
software directly on our OpenCL kernels, you should consider waiting for this
update. You can follow development live on the
[development branch](https://github.com/clij/clij-opencl-kernels/tree/development/src/main/java/net/haesleinhuepf/clij/kernels).
6 changes: 4 additions & 2 deletions license.txt
View file
Original file line number Diff line number Diff line change
@@ -1,5 +1,7 @@
Copyright 2019 Robert Haase, Nico Stuurman, Max Planck Institute for Molecular Cell Biology
and Genetics Dresden, Regents of the Univeristy of California
Copyright 2019 Robert Haase, Nico Stuurman, Deborah Schmidt, Uwe Schmidt, Martin Weigert, Peter Haub, Fabrice P. Cordeli�res
Max Planck Institute for Molecular Cell Biology
and Genetics Dresden, University of Virginia,
Regents of the Univeristy of California

Redistribution and use in source and binary forms, with or without
modification, are permitted provided that the following conditions are met:
Expand Down
14 changes: 11 additions & 3 deletions pom.xml
View file
Original file line number Diff line number Diff line change
Expand Up @@ -8,13 +8,13 @@
<parent>
<groupId>net.haesleinhuepf</groupId>
<artifactId>clij-parent-pom</artifactId>
<version>1.4.1</version>
<version>2.0.0.10</version>
<relativePath />
</parent>

<groupId>net.haesleinhuepf</groupId>
<artifactId>clij-opencl-kernels</artifactId>
<version>1.4.1</version>
<version>2.0.0.10</version>

<name>clij-opencl-kernels</name>
<description>clij-opencl-kernels</description>
Expand Down Expand Up @@ -49,9 +49,17 @@
</developer>
</developers>
<contributors>
<contributor>
<contributor>
<name>Nico Stuurman</name>
</contributor>
<contributor>
<name>Peter Haub</name>
<properties><id>iwbh15</id></properties>
</contributor>
<contributor>
<name>Ruth Whelan-Jeans</name>
<properties><id>ruthwj</id></properties>
</contributor>
</contributors>

<mailingLists>
Expand Down
8 changes: 4 additions & 4 deletions ...haesleinhuepf/clij/kernels/absolute_2d.cl → ...esleinhuepf/clij/kernels/absolute_2d_x.cl
View file
Original file line number Diff line number Diff line change
@@ -1,18 +1,18 @@
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

__kernel void absolute_2d(DTYPE_IMAGE_IN_2D src,
DTYPE_IMAGE_OUT_2D dst
__kernel void absolute_2d(IMAGE_src_TYPE src,
IMAGE_dst_TYPE dst
)
{
const int x = get_global_id(0);
const int y = get_global_id(1);

const int2 pos = (int2){x,y};

float value = READ_IMAGE_2D(src, sampler, pos).x;
float value = READ_src_IMAGE(src, sampler, pos).x;
if ( value < 0 ) {
value = -1 * value;
}

WRITE_IMAGE_2D (dst, pos, CONVERT_DTYPE_OUT(value));
WRITE_dst_IMAGE (dst, pos, CONVERT_dst_PIXEL_TYPE(value));
}
8 changes: 4 additions & 4 deletions ...haesleinhuepf/clij/kernels/absolute_3d.cl → ...esleinhuepf/clij/kernels/absolute_3d_x.cl
View file
Original file line number Diff line number Diff line change
@@ -1,7 +1,7 @@
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

__kernel void absolute_3d(DTYPE_IMAGE_IN_3D src,
DTYPE_IMAGE_OUT_3D dst
__kernel void absolute_3d(IMAGE_src_TYPE src,
IMAGE_dst_TYPE dst
)
{
const int x = get_global_id(0);
Expand All @@ -10,10 +10,10 @@ __kernel void absolute_3d(DTYPE_IMAGE_IN_3D src,

const int4 pos = (int4){x,y,z,0};

float value = READ_IMAGE_3D(src, sampler, pos).x;
float value = READ_src_IMAGE(src, sampler, pos).x;
if ( value < 0 ) {
value = -1 * value;
}

WRITE_IMAGE_3D (dst, pos, CONVERT_DTYPE_OUT(value));
WRITE_dst_IMAGE (dst, pos, CONVERT_dst_PIXEL_TYPE(value));
}
19 changes: 0 additions & 19 deletions src/main/java/net/haesleinhuepf/clij/kernels/addWeightedPixelwise_2d.cl
View file

This file was deleted.

19 changes: 0 additions & 19 deletions src/main/java/net/haesleinhuepf/clij/kernels/addWeightedPixelwise_3d.cl
View file

This file was deleted.

17 changes: 17 additions & 0 deletions src/main/java/net/haesleinhuepf/clij/kernels/add_image_and_scalar_2d_x.cl
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,17 @@
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

__kernel void add_image_and_scalar_2d(
IMAGE_src_TYPE src,
IMAGE_dst_TYPE dst,
float scalar
)
{
const int x = get_global_id(0);
const int y = get_global_id(1);

const int2 pos = (int2){x,y};

const IMAGE_dst_PIXEL_TYPE value = CONVERT_dst_PIXEL_TYPE(READ_src_IMAGE(src, sampler, pos).x + scalar);

WRITE_dst_IMAGE (dst, pos, value);
}
18 changes: 18 additions & 0 deletions src/main/java/net/haesleinhuepf/clij/kernels/add_image_and_scalar_3d_x.cl
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,18 @@
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

__kernel void add_image_and_scalar_3d(
IMAGE_src_TYPE src,
IMAGE_dst_TYPE dst,
float scalar
)
{
const int x = get_global_id(0);
const int y = get_global_id(1);
const int z = get_global_id(2);

const int4 pos = (int4){x,y,z,0};

const IMAGE_dst_PIXEL_TYPE value = CONVERT_dst_PIXEL_TYPE(READ_src_IMAGE(src, sampler, pos).x + scalar);

WRITE_dst_IMAGE(dst, pos, value);
}
Loading

0 comments on commit 9546d0f

Please sign in to comment.