Good call. I have to explicitly load gcc, and I didn’t think this would need it. Unfortunately, after loading gcc/8.3.0, I get a different error… This one is a bit long, so I apologize in advance.
Console Output
package net.haesleinhuepf.clijx.plugins, clij2-fft, version 0.0
//###########################################################################
// Preamble:
#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable
#pragma OPENCL EXTENSION cl_amd_printf : enable
#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
#ifndef M_PI
#define M_PI 3.14159265358979323846f /* pi */
#endif
#ifndef M_LOG2E
#define M_LOG2E 1.4426950408889634074f /* log_2 e */
#endif
#ifndef M_LOG10E
#define M_LOG10E 0.43429448190325182765f /* log_10 e */
#endif
#ifndef M_LN2
#define M_LN2 0.69314718055994530942f /* log_e 2 */
#endif
#ifndef M_LN10
#define M_LN10 2.30258509299404568402f /* log_e 10 */
#endif
#ifndef BUFFER_READ_WRITE
#define BUFFER_READ_WRITE 1
#define MINMAX_TYPE long
inline char2 read_buffer4dc(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global char * buffer_var, sampler_t sampler, int4 position )
{
int4 pos = (int4){position.x, position.y, position.z, position.w};
if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
// todo: correct pos.w
}
long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height + pos.w * read_buffer_width * read_buffer_height * read_buffer_depth ;
if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) { // todo: check pos.w
return (char2){0, 0};
}
return (char2){buffer_var[pos_in_buffer],0};
}
inline uchar2 read_buffer4duc(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global uchar * buffer_var, sampler_t sampler, int4 position )
{
int4 pos = (int4){position.x, position.y, position.z, position.w};
if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
// todo: correct pos.w
}
long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height + pos.w * read_buffer_width * read_buffer_height * read_buffer_depth;
if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) { // todo: check pos.w
return (uchar2){0, 0};
}
return (uchar2){buffer_var[pos_in_buffer],0};
}
inline short2 read_buffer4di(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global short * buffer_var, sampler_t sampler, int4 position )
{
int4 pos = (int4){position.x, position.y, position.z, position.w};
if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
// todo: correct pos.w
}
long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height + pos.w * read_buffer_width * read_buffer_height * read_buffer_depth;
if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) { // todo: check pos.w
return (short2){0, 0};
}
return (short2){buffer_var[pos_in_buffer],0};
}
inline ushort2 read_buffer4dui(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global ushort * buffer_var, sampler_t sampler, int4 position )
{
int4 pos = (int4){position.x, position.y, position.z, position.w};
if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
// todo: correct pos.w
}
long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height + pos.w * read_buffer_width * read_buffer_height * read_buffer_depth;
if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) { // todo: check pos.w
return (ushort2){0, 0};
}
return (ushort2){buffer_var[pos_in_buffer],0};
}
inline float2 read_buffer4df(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global float* buffer_var, sampler_t sampler, int4 position )
{
int4 pos = (int4){position.x, position.y, position.z, position.w};
if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
// todo: correct pos.w
}
long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height + pos.w * read_buffer_width * read_buffer_height * read_buffer_depth;
if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) { // todo: check pos.w
return (float2){0, 0};
}
return (float2){buffer_var[pos_in_buffer],0};
}
inline void write_buffer4dc(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global char * buffer_var, int4 pos, char value )
{
long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width + (long)pos.z * write_buffer_width * write_buffer_height + pos.w * write_buffer_width * write_buffer_height * write_buffer_depth;
if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height || pos.z < 0 || (long)pos.z >= write_buffer_depth) { // todo: check pos.w
return;
}
buffer_var[pos_in_buffer] = value;
}
inline void write_buffer4duc(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global uchar * buffer_var, int4 pos, uchar value )
{
long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width + (long)pos.z * write_buffer_width * write_buffer_height + pos.w * write_buffer_width * write_buffer_height * write_buffer_depth;
if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height || pos.z < 0 || (long)pos.z >= write_buffer_depth) { // todo: check pos.w
return;
}
buffer_var[pos_in_buffer] = value;
}
inline void write_buffer4di(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global short * buffer_var, int4 pos, short value )
{
long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width + (long)pos.z * write_buffer_width * write_buffer_height + pos.w * write_buffer_width * write_buffer_height * write_buffer_depth;
if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height || pos.z < 0 || (long)pos.z >= write_buffer_depth) { // todo: check pos.w
return;
}
buffer_var[pos_in_buffer] = value;
}
inline void write_buffer4dui(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global ushort * buffer_var, int4 pos, ushort value )
{
long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width + (long)pos.z * write_buffer_width * write_buffer_height + pos.w * write_buffer_width * write_buffer_height * write_buffer_depth;
if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height || pos.z < 0 || (long)pos.z >= write_buffer_depth) { // todo: check pos.w
return;
}
buffer_var[pos_in_buffer] = value;
}
inline void write_buffer4df(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global float* buffer_var, int4 pos, float value )
{
long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width + (long)pos.z * write_buffer_width * write_buffer_height + pos.w * write_buffer_width * write_buffer_height * write_buffer_depth;
if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height || pos.z < 0 || (long)pos.z >= write_buffer_depth) { // todo: check pos.w
return;
}
buffer_var[pos_in_buffer] = value;
}
inline char2 read_buffer3dc(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global char * buffer_var, sampler_t sampler, int4 position )
{
int4 pos = (int4){position.x, position.y, position.z, 0};
if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
}
long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height;
if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) {
return (char2){0, 0};
}
return (char2){buffer_var[pos_in_buffer],0};
}
inline uchar2 read_buffer3duc(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global uchar * buffer_var, sampler_t sampler, int4 position )
{
int4 pos = (int4){position.x, position.y, position.z, 0};
if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
}
long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height;
if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) {
return (uchar2){0, 0};
}
return (uchar2){buffer_var[pos_in_buffer],0};
}
inline short2 read_buffer3di(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global short * buffer_var, sampler_t sampler, int4 position )
{
int4 pos = (int4){position.x, position.y, position.z, 0};
if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
}
long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height;
if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) {
return (short2){0, 0};
}
return (short2){buffer_var[pos_in_buffer],0};
}
inline ushort2 read_buffer3dui(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global ushort * buffer_var, sampler_t sampler, int4 position )
{
int4 pos = (int4){position.x, position.y, position.z, 0};
if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
}
long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height;
if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) {
return (ushort2){0, 0};
}
return (ushort2){buffer_var[pos_in_buffer],0};
}
inline float2 read_buffer3df(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global float* buffer_var, sampler_t sampler, int4 position )
{
int4 pos = (int4){position.x, position.y, position.z, 0};
if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
}
long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height;
if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) {
return (float2){0, 0};
}
return (float2){buffer_var[pos_in_buffer],0};
}
inline void write_buffer3dc(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global char * buffer_var, int4 pos, char value )
{
long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width + (long)pos.z * write_buffer_width * write_buffer_height;
if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height || pos.z < 0 || (long)pos.z >= write_buffer_depth) {
return;
}
buffer_var[pos_in_buffer] = value;
}
inline void write_buffer3duc(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global uchar * buffer_var, int4 pos, uchar value )
{
long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width + (long)pos.z * write_buffer_width * write_buffer_height;
if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height || pos.z < 0 || (long)pos.z >= write_buffer_depth) {
return;
}
buffer_var[pos_in_buffer] = value;
}
inline void write_buffer3di(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global short * buffer_var, int4 pos, short value )
{
long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width + (long)pos.z * write_buffer_width * write_buffer_height;
if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height || pos.z < 0 || (long)pos.z >= write_buffer_depth) {
return;
}
buffer_var[pos_in_buffer] = value;
}
inline void write_buffer3dui(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global ushort * buffer_var, int4 pos, ushort value )
{
long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width + (long)pos.z * write_buffer_width * write_buffer_height;
if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height || pos.z < 0 || (long)pos.z >= write_buffer_depth) {
return;
}
buffer_var[pos_in_buffer] = value;
}
inline void write_buffer3df(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global float* buffer_var, int4 pos, float value )
{
long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width + (long)pos.z * write_buffer_width * write_buffer_height;
if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height || pos.z < 0 || (long)pos.z >= write_buffer_depth) {
return;
}
buffer_var[pos_in_buffer] = value;
}
inline char2 read_buffer2dc(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global char * buffer_var, sampler_t sampler, int2 position )
{
int2 pos = (int2){position.x, position.y};
if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
}
long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width;
if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height) {
return (char2){0, 0};
}
return (char2){buffer_var[pos_in_buffer],0};
}
inline uchar2 read_buffer2duc(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global uchar * buffer_var, sampler_t sampler, int2 position )
{
int2 pos = (int2){position.x, position.y};
if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
}
long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width;
if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height) {
return (uchar2){0, 0};
}
return (uchar2){buffer_var[pos_in_buffer],0};
}
inline short2 read_buffer2di(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global short * buffer_var, sampler_t sampler, int2 position )
{
int2 pos = (int2){position.x, position.y};
if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
}
long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width;
if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height) {
return (short2){0, 0};
}
return (short2){buffer_var[pos_in_buffer],0};
}
inline ushort2 read_buffer2dui(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global ushort * buffer_var, sampler_t sampler, int2 position )
{
int2 pos = (int2){position.x, position.y};
if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
}
long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width;
if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height) {
return (ushort2){0, 0};
}
return (ushort2){buffer_var[pos_in_buffer],0};
}
inline float2 read_buffer2df(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global float* buffer_var, sampler_t sampler, int2 position )
{
int2 pos = (int2){position.x, position.y};
if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
}
long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width;
if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height) {
return (float2){0, 0};
}
return (float2){buffer_var[pos_in_buffer],0};
}
inline void write_buffer2dc(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global char * buffer_var, int2 pos, char value )
{
long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width;
if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height) {
return;
}
buffer_var[pos_in_buffer] = value;
}
inline void write_buffer2duc(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global uchar * buffer_var, int2 pos, uchar value )
{
long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width;
if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height) {
return;
}
buffer_var[pos_in_buffer] = value;
}
inline void write_buffer2di(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global short * buffer_var, int2 pos, short value )
{
long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width;
if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height) {
return;
}
buffer_var[pos_in_buffer] = value;
}
inline void write_buffer2dui(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global ushort * buffer_var, int2 pos, ushort value )
{
long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width;
if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height) {
return;
}
buffer_var[pos_in_buffer] = value;
}
inline void write_buffer2df(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global float* buffer_var, int2 pos, float value )
{
long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width;
if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height) {
return;
}
buffer_var[pos_in_buffer] = value;
}
inline uchar clij_convert_uchar_sat(float value) {
if (value > 255) {
return 255;
}
if (value < 0) {
return 0;
}
return (uchar)value;
}
inline char clij_convert_char_sat(float value) {
if (value > 127) {
return 127;
}
if (value < -128) {
return -128;
}
return (char)value;
}
inline ushort clij_convert_ushort_sat(float value) {
if (value > 65535) {
return 65535;
}
if (value < 0) {
return 0;
}
return (ushort)value;
}
inline short clij_convert_short_sat(float value) {
if (value > 32767) {
return 32767;
}
if (value < -32768) {
return -32768;
}
return (short)value;
}
inline uint clij_convert_uint_sat(float value) {
if (value > 4294967295) {
return 4294967295;
}
if (value < 0) {
return 0;
}
return (uint)value;
}
inline int clij_convert_int_sat(float value) {
if (value > 2147483647) {
return 2147483647;
}
if (value < -2147483648) {
return -2147483648;
}
return (int)value;
}
inline float clij_convert_float_sat(float value) {
return value;
}
#define READ_IMAGE(a,b,c) READ_ ## a ## IMAGE(a,b,c)
#define WRITE_IMAGE(a,b,c) WRITE ## a ## _IMAGE(a,b,c)
#endif
//###########################################################################
// Defines:
#define WRITE_dst_IMAGE(a,b,c) write_buffer3df(GET_IMAGE_WIDTH(a),GET_IMAGE_HEIGHT(a),GET_IMAGE_DEPTH(a),a,b,c)
#define IMAGE_src_PIXEL_TYPE ushort
#define CONVERT_dst_PIXEL_TYPE clij_convert_float_sat
#define POS_src_INSTANCE(pos0, pos1, pos2, pos3) ((int4)(pos0, pos1, pos2, pos3))
#define GET_IMAGE_HEIGHT(image_key) image_size_ ## image_key ## height
#define READ_src_IMAGE(a,b,c) read_buffer3dui(GET_IMAGE_WIDTH(a),GET_IMAGE_HEIGHT(a),GET_IMAGE_DEPTH(a),a,b,c)
#define POS_dst_INSTANCE(pos0, pos1, pos2, pos3) ((int4)(pos0, pos1, pos2, pos3))
#define POS_src_TYPE int4
#define GET_IMAGE_DEPTH(image_key) image_size ## image_key ## _depth
#define IMAGE_dst_TYPE long image_size_dst_width, long image_size_dst_height, long image_size_dst_depth, _global float*
#define MAX_ARRAY_SIZE 1000
#define CONVERT_src_PIXEL_TYPE clij_convert_ushort_sat
#define WRITE_src_IMAGE(a,b,c) write_buffer3dui(GET_IMAGE_WIDTH(a),GET_IMAGE_HEIGHT(a),GET_IMAGE_DEPTH(a),a,b,c)
#define IMAGE_dst_PIXEL_TYPE float
#define GET_IMAGE_WIDTH(image_key) image_size ## image_key ## _width
#define IMAGE_src_TYPE long image_size_src_width, long image_size_src_height, long image_size_src_depth, __global ushort*
#define READ_dst_IMAGE(a,b,c) read_buffer3df(GET_IMAGE_WIDTH(a),GET_IMAGE_HEIGHT(a),GET_IMAGE_DEPTH(a),a,b,c)
#define POS_dst_TYPE int4
//###########################################################################
// Source: ‘crop_3d_x.cl’ relative to Crop3D
__kernel void crop_3d(
IMAGE_dst_TYPE dst,
IMAGE_src_TYPE src,
int start_x,
int start_y,
int start_z
) {
const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
const int dx = get_global_id(0);
const int dy = get_global_id(1);
const int dz = get_global_id(2);
const int sx = start_x + dx;
const int sy = start_y + dy;
const int sz = start_z + dz;
const POS_dst_TYPE dpos = POS_dst_INSTANCE(dx, dy, dz, 0);
const POS_src_TYPE spos = POS_src_INSTANCE(sx, sy, sz, 0);
const float out = READ_src_IMAGE(src,sampler,spos).x;
WRITE_dst_IMAGE(dst,dpos, CONVERT_dst_PIXEL_TYPE(out));
}
Error when trying to create kernel crop_3d
net.haesleinhuepf.clij.clearcl.exceptions.OpenCLException: OpenCL error: -45 → CL_INVALID_PROGRAM_EXECUTABLE
at net.haesleinhuepf.clij.clearcl.backend.BackendUtils.checkOpenCLErrorCode(BackendUtils.java:352)
at net.haesleinhuepf.clij.clearcl.backend.jocl.ClearCLBackendJOCL.lambda$getKernelPeerPointer$19(ClearCLBackendJOCL.java:601)
at net.haesleinhuepf.clij.clearcl.backend.BackendUtils.checkExceptions(BackendUtils.java:156)
at net.haesleinhuepf.clij.clearcl.backend.jocl.ClearCLBackendJOCL.getKernelPeerPointer(ClearCLBackendJOCL.java:593)
at net.haesleinhuepf.clij.clearcl.ClearCLCompiledProgram.createKernel(ClearCLCompiledProgram.java:137)
at net.haesleinhuepf.clij.clearcl.ClearCLProgram.createKernel(ClearCLProgram.java:685)
at net.haesleinhuepf.clij.clearcl.util.CLKernelExecutor.getKernel(CLKernelExecutor.java:353)
at net.haesleinhuepf.clij.clearcl.util.CLKernelExecutor.enqueue(CLKernelExecutor.java:229)
at net.haesleinhuepf.clij2.CLIJ2.lambda$executeSubsequently$0(CLIJ2.java:466)
at net.haesleinhuepf.clij.clearcl.util.ElapsedTime.measure(ElapsedTime.java:97)
at net.haesleinhuepf.clij.clearcl.util.ElapsedTime.measure(ElapsedTime.java:28)
at net.haesleinhuepf.clij2.CLIJ2.executeSubsequently(CLIJ2.java:456)
at net.haesleinhuepf.clij2.CLIJ2.executeSubsequently(CLIJ2.java:443)
at net.haesleinhuepf.clij2.CLIJ2.executeSubsequently(CLIJ2.java:438)
at net.haesleinhuepf.clij2.CLIJ2.execute(CLIJ2.java:423)
at net.haesleinhuepf.clij2.plugins.Crop3D.crop(Crop3D.java:66)
at net.haesleinhuepf.clij2.CLIJ2Ops.crop(CLIJ2Ops.java:3624)
at net.haesleinhuepf.clijx.plugins.OpenCLFFTUtility.cropExtended(OpenCLFFTUtility.java:260)
at net.haesleinhuepf.clijx.plugins.DeconvolveRichardsonLucyFFT.deconvolveFFT(DeconvolveRichardsonLucyFFT.java:112)
at net.haesleinhuepf.clijx.plugins.DeconvolveRichardsonLucyFFT.deconvolveRichardsonLucyFFT(DeconvolveRichardsonLucyFFT.java:77)
at net.haesleinhuepf.clijx.plugins.DeconvolveRichardsonLucyFFT.executeCL(DeconvolveRichardsonLucyFFT.java:42)
at net.haesleinhuepf.clij.macro.AbstractCLIJPlugin.run(AbstractCLIJPlugin.java:478)
at ij.plugin.filter.PlugInFilterRunner.processOneImage(PlugInFilterRunner.java:265)
at ij.plugin.filter.PlugInFilterRunner.(PlugInFilterRunner.java:114)
at ij.IJ.runUserPlugIn(IJ.java:237)
at ij.IJ.runPlugIn(IJ.java:198)
at ij.Executer.runCommand(Executer.java:150)
at ij.Executer.run(Executer.java:68)
at java.lang.Thread.run(Thread.java:748)
