opencl.c

Go to the documentation of this file.

/*
    This file is part of darktable,
    Copyright (C) 2010-2012, 2016 johannes hanika.
    Copyright (C) 2011 Bruce Guenter.
    Copyright (C) 2011 Henrik Andersson.
    Copyright (C) 2011 Moritz Lipp.
    Copyright (C) 2011-2019 Ulrich Pegelow.
    Copyright (C) 2012 Christian Tellefsen.
    Copyright (C) 2012 Jérémy Rosen.
    Copyright (C) 2012 Michal Babej.
    Copyright (C) 2012 Richard Wonka.
    Copyright (C) 2012-2014, 2016-2018 Tobias Ellinghaus.
    Copyright (C) 2013-2019 Roman Lebedev.
    Copyright (C) 2013 Simon Spannagel.
    Copyright (C) 2015, 2019 Dan Torop.
    Copyright (C) 2015, 2017 parafin.
    Copyright (C) 2015 Pascal de Bruijn.
    Copyright (C) 2016-2017, 2019 Peter Budai.
    Copyright (C) 2017-2019 Edgardo Hoszowski.
    Copyright (C) 2017, 2019 luzpaz.
    Copyright (C) 2019 Andreas Schneider.
    Copyright (C) 2019, 2021-2022, 2025-2026 Aurélien PIERRE.
    Copyright (C) 2019 Damian D. Martinez Dreyer.
    Copyright (C) 2019-2020 Heiko Bauke.
    Copyright (C) 2019 jakubfi.
    Copyright (C) 2019-2021 Pascal Obry.
    Copyright (C) 2020 David-Tillmann Schaefer.
    Copyright (C) 2020-2022 Hubert Kowalski.
    Copyright (C) 2020-2021 Ralf Brown.
    Copyright (C) 2021 Chris Elston.
    Copyright (C) 2022 Hanno Schwalm.
    Copyright (C) 2022 Martin Bařinka.
    Copyright (C) 2022 Victor Forsiuk.
    Copyright (C) 2024 Alynx Zhou.
    Copyright (C) 2025 Guillaume Stutin.
    
    darktable is free software: you can redistribute it and/or modify
    it under the terms of the GNU General Public License as published by
    the Free Software Foundation, either version 3 of the License, or
    (at your option) any later version.
    
    darktable is distributed in the hope that it will be useful,
    but WITHOUT ANY WARRANTY; without even the implied warranty of
    MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
    GNU General Public License for more details.
    
    You should have received a copy of the GNU General Public License
    along with darktable.  If not, see <http://www.gnu.org/licenses/>.
*/
 
#ifdef HAVE_OPENCL
 
#include "common/opencl.h"
#include "common/bilateralcl.h"
#include "common/darktable.h"
#include "common/dlopencl.h"
#include "common/dwt.h"
#include "common/file_location.h"
#include "common/gaussian.h"
#include "common/guided_filter.h"
#include "common/heal.h"
#include "common/interpolation.h"
#include "common/locallaplaciancl.h"
#include "common/nvidia_gpus.h"
#include "common/opencl_drivers_blacklist.h"
#include "common/tea.h"
#include "control/conf.h"
#include "control/control.h"
#include "gui/splash.h"
#include "develop/blend.h"
#include "develop/pixelpipe.h"
#include "develop/pixelpipe_cache.h"
 
#include <assert.h>
#include <locale.h>
#include <stdio.h>
#include <string.h>
#include <strings.h>
 
#include <ctype.h>
#include <errno.h>
#include <libgen.h>
#include <sys/stat.h>
#include <zlib.h>
 
static gboolean _opencl_splash_active = FALSE;
 
static inline void _opencl_splash_update_compile(const char *programname)
{
  if(IS_NULL_PTR(programname)) return;
  if(IS_NULL_PTR(darktable.gui)) return;
 
  if(!_opencl_splash_active)
  {
    dt_gui_splash_init();
    _opencl_splash_active = TRUE;
  }
 
  dt_gui_splash_updatef(_("Building OpenCL kernels: %s"), programname);
}
 
static const char *dt_opencl_get_vendor_by_id(unsigned int id);
static char *_ascii_str_canonical(const char *in, char *out, int maxlen);
static void dt_opencl_priority_parse(dt_opencl_t *cl, char *configstr, int *priority_list, int *mandatory);
static void dt_opencl_update_priorities();
static void dt_opencl_apply_scheduling_profile();
static void dt_opencl_set_synchronization_timeout(int value);
 
 
int dt_opencl_get_device_info(dt_opencl_t *cl, cl_device_id device, cl_device_info param_name, void **param_value,
                              size_t *param_value_size)
{
  *param_value_size = SIZE_MAX;
 
  // 1. figure out how much memory is needed
  cl_int err = (cl->dlocl->symbols->dt_clGetDeviceInfo)(device, param_name, 0, NULL, param_value_size);
  if(err != CL_SUCCESS)
  {
    dt_print(DT_DEBUG_OPENCL,
             "[dt_opencl_get_device_info] could not query the actual size in bytes of info %d: %i\n", param_name, err);
    goto error;
  }
 
  // 2. did we /actually/ get the size?
  if(*param_value_size == SIZE_MAX || *param_value_size == 0)
  {
    // both of these sizes make no sense. either i failed to parse spec, or opencl implementation bug?
    dt_print(DT_DEBUG_OPENCL,
             "[dt_opencl_get_device_info] ERROR: no size returned, or zero size returned for data %d: %" G_GSIZE_FORMAT "\n",
             param_name, *param_value_size);
    err = CL_INVALID_VALUE; // FIXME: anything better?
    goto error;
  }
 
  // 3. make sure that *param_value points to big-enough memory block
  {
    void *ptr = realloc(*param_value, *param_value_size);
    if(IS_NULL_PTR(ptr))
    {
      dt_print(DT_DEBUG_OPENCL,
               "[dt_opencl_get_device_info] memory allocation failed! tried to allocate %" G_GSIZE_FORMAT " bytes for data %d: %i",
               *param_value_size, param_name, err);
      err = CL_OUT_OF_HOST_MEMORY;
      goto error;
    }
 
    // allocation succeeded, update pointer.
    *param_value = ptr;
  }
 
  // 4. actually get the value
  err = (cl->dlocl->symbols->dt_clGetDeviceInfo)(device, param_name, *param_value_size, *param_value, NULL);
  if(err != CL_SUCCESS)
  {
    dt_print(DT_DEBUG_OPENCL, "[dt_opencl_get_device_info] could not query info %d: %i\n", param_name, err);
    goto error;
  }
 
  return CL_SUCCESS;
 
error:
  dt_free(*param_value);
  *param_value_size = 0;
  return err;
}
 
int dt_opencl_avoid_atomics(const int devid)
{
  dt_opencl_t *cl = darktable.opencl;
  return (!cl->inited || devid < 0) ? 0 : cl->dev[devid].avoid_atomics;
}
 
int dt_opencl_micro_nap(const int devid)
{
  dt_opencl_t *cl = darktable.opencl;
  return (!cl->inited || devid < 0) ? 0 : cl->dev[devid].micro_nap;
}
 
gboolean dt_opencl_use_pinned_memory(const int devid)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited || devid < 0) return FALSE;
  return cl->dev[devid].pinned_memory & DT_OPENCL_PINNING_ON;
}
 
gboolean dt_opencl_is_pinned_memory(cl_mem mem)
{
  const cl_mem_flags flags = dt_opencl_get_mem_flags(mem);
  return (flags & CL_MEM_USE_HOST_PTR) || (flags & CL_MEM_ALLOC_HOST_PTR);
}
 
void dt_opencl_write_device_config(const int devid)
{
  if(devid < 0) return;
  dt_opencl_t *cl = darktable.opencl;
  gchar buf[256] = { 0 };
  gchar key_device[256] = { 0 };
  g_snprintf(key_device, 254, "%s/%i/%s", DT_CLDEVICE_HEAD, devid, cl->dev[devid].cname);
 
  g_snprintf(buf, sizeof(buf), "%s/avoid_atomics", key_device);
  dt_conf_set_int(buf, cl->dev[devid].avoid_atomics);
 
  g_snprintf(buf, sizeof(buf), "%s/micro_nap", key_device);
  dt_conf_set_int(buf, cl->dev[devid].micro_nap);
 
  g_snprintf(buf, sizeof(buf), "%s/pinned_memory", key_device);
  dt_conf_set_int(buf, cl->dev[devid].pinned_memory & (DT_OPENCL_PINNING_ON | DT_OPENCL_PINNING_DISABLED));
 
  g_snprintf(buf, sizeof(buf), "%s/wd", key_device);
  dt_conf_set_int(buf, cl->dev[devid].clroundup_wd);
 
  g_snprintf(buf, sizeof(buf), "%s/ht", key_device);
  dt_conf_set_int(buf, cl->dev[devid].clroundup_ht);
 
  g_snprintf(buf, sizeof(buf), "%s/event_handles", key_device);
  dt_conf_set_int(buf, cl->dev[devid].event_handles);
 
  g_snprintf(buf, sizeof(buf), "%s/disabled", key_device);
  dt_conf_set_int(buf, cl->dev[devid].disabled & 1);
 
  g_snprintf(buf, sizeof(buf), "%s/id%i/forced_headroom", key_device, devid);
  dt_conf_set_int(buf, cl->dev[devid].forced_headroom);
}
 
static int _dt_opencl_get_conf_int(const gchar *key_device, const gchar *conf_name, gboolean *safety_ok)
{
  int res = 0;
  gchar *key = g_strconcat(key_device, "/", conf_name, NULL);
  const gboolean existing_device = dt_conf_key_not_empty(key);
  if(existing_device)
    res = dt_conf_get_int(key);
  else
  {
    dt_print(DT_DEBUG_OPENCL, "Warning: conf '%s' not found in anselrc.\n", key);
    *safety_ok = FALSE;
  }
 
  dt_free(key);
  return res;
}
 
gboolean dt_opencl_read_device_config(const int devid)
{
  if(devid < 0) return FALSE;
  dt_opencl_t *cl = darktable.opencl;
  gchar key_device[256] = { 0 };
  g_snprintf(key_device, 254, "%s/%i/%s", DT_CLDEVICE_HEAD, devid, cl->dev[devid].cname);
  gboolean safety_ok = TRUE;
 
  int avoid_atomics = _dt_opencl_get_conf_int(key_device, "avoid_atomics", &safety_ok);
  int micro_nap = _dt_opencl_get_conf_int(key_device, "micro_nap", &safety_ok);
  int pinned_memory = _dt_opencl_get_conf_int(key_device, "pinned_memory", &safety_ok);
  int wd = _dt_opencl_get_conf_int(key_device, "wd", &safety_ok);
  int ht = _dt_opencl_get_conf_int(key_device, "ht", &safety_ok);
  int event_handles = _dt_opencl_get_conf_int(key_device, "event_handles", &safety_ok);
  int disabled = _dt_opencl_get_conf_int(key_device, "disabled", &safety_ok);
 
  // some rudimentary safety checking if string seems to be ok
  safety_ok |= (wd > 1) && (wd < 513) && (ht > 1) && (ht < 513);
 
  if(safety_ok)
  {
    cl->dev[devid].avoid_atomics = avoid_atomics;
    cl->dev[devid].micro_nap = micro_nap;
    cl->dev[devid].pinned_memory = pinned_memory;
    cl->dev[devid].clroundup_wd = wd;
    cl->dev[devid].clroundup_ht = ht;
    cl->dev[devid].event_handles = event_handles;
    cl->dev[devid].disabled = disabled;
  }
  else // if there is something wrong with the found conf key reset to defaults
  {
    dt_print(DT_DEBUG_OPENCL, "[dt_opencl_read_device_config] malformed data '%s'\n", key_device);
  }
 
  // do some safety housekeeping
  cl->dev[devid].avoid_atomics &= 1;
  cl->dev[devid].pinned_memory &= (DT_OPENCL_PINNING_ON | DT_OPENCL_PINNING_DISABLED);
  cl->dev[devid].micro_nap = CLAMP(cl->dev[devid].micro_nap, 250, 1000000);
  if((cl->dev[devid].clroundup_wd < 2) || (cl->dev[devid].clroundup_wd > 512))
    cl->dev[devid].clroundup_wd = 16;
  if((cl->dev[devid].clroundup_ht < 2) || (cl->dev[devid].clroundup_ht > 512))
    cl->dev[devid].clroundup_ht = 16;
  if(cl->dev[devid].event_handles < 0)
    cl->dev[devid].event_handles = 0x40961440;
 
  cl->dev[devid].use_events = (cl->dev[devid].event_handles != 0) ? 1 : 0;
  cl->dev[devid].disabled &= 1;
 
  // Also take care of extended device data, these are not only device specific but also depend on the devid
  g_snprintf(key_device, 254, "%s/%i/%s/id%i/forced_headroom", DT_CLDEVICE_HEAD, devid, cl->dev[devid].cname, devid);
  if(dt_conf_key_not_empty(key_device))
  {
    int forced_headroom = dt_conf_get_int(key_device);
    if(forced_headroom > 0) cl->dev[devid].forced_headroom = forced_headroom;
  }
  else if(cl->dev[devid].host_unified_memory)
    cl->dev[devid].forced_headroom = MAX(512ul, (size_t)dt_conf_get_int64("memory_opencl_headroom"));
  else // this is used if updating to 4.0 or fresh installs; see commenting _opencl_get_unused_device_mem()
    cl->dev[devid].forced_headroom = dt_conf_get_int64("memory_opencl_headroom");
 
  dt_opencl_write_device_config(devid);
  return !safety_ok;
}
 
int dt_opencl_get_detected_device_count(void)
{
  dt_opencl_t *cl = darktable.opencl;
  if(IS_NULL_PTR(cl)) return 0;
 
  return cl->num_detected_devs;
}
 
const dt_opencl_detected_device_t *dt_opencl_get_detected_device(const int detected)
{
  dt_opencl_t *cl = darktable.opencl;
  if(IS_NULL_PTR(cl) || detected < 0 || detected >= cl->num_detected_devs) return NULL;
 
  return cl->detected_devs + detected;
}
 
gboolean dt_opencl_detected_device_enabled(const int detected)
{
  const dt_opencl_detected_device_t *device = dt_opencl_get_detected_device(detected);
  if(IS_NULL_PTR(device)) return FALSE;
 
  gchar key[256] = { 0 };
  g_snprintf(key, sizeof(key), "%s/%d/%s/disabled", DT_CLDEVICE_HEAD, device->config_id,
             !IS_NULL_PTR(device->cname) ? device->cname : "");
  const gboolean disabled = dt_conf_key_not_empty(key) ? dt_conf_get_int(key) : (device->disabled & 1);
 
  return !disabled;
}
 
int dt_opencl_set_detected_device_enabled(const int detected, const gboolean enabled)
{
  const dt_opencl_detected_device_t *device = dt_opencl_get_detected_device(detected);
  if(IS_NULL_PTR(device)) return -1;
 
  gchar key[256] = { 0 };
  g_snprintf(key, sizeof(key), "%s/%d/%s/disabled", DT_CLDEVICE_HEAD, device->config_id,
             !IS_NULL_PTR(device->cname) ? device->cname : "");
  dt_conf_set_int(key, enabled ? 0 : 1);
 
  dt_opencl_t *cl = darktable.opencl;
  cl->detected_devs[detected].disabled = enabled ? 0 : 1;
 
  gboolean opencl_enabled = enabled;
  if(!opencl_enabled)
  {
    // The global OpenCL preference is derived from all detected GPUs. We are looking
    // for any GPU still enabled before turning OpenCL off globally.
    for(int dev = 0; dev < cl->num_detected_devs; dev++)
    {
      if(dt_opencl_detected_device_enabled(dev))
      {
        opencl_enabled = TRUE;
        break;
      }
    }
  }
 
  dt_conf_set_bool("opencl", opencl_enabled);
  return 0;
}
 
gboolean dt_opencl_detected_device_pinned_memory(const int detected)
{
  const dt_opencl_detected_device_t *device = dt_opencl_get_detected_device(detected);
  if(IS_NULL_PTR(device)) return FALSE;
 
  gchar key[256] = { 0 };
  g_snprintf(key, sizeof(key), "%s/%d/%s/pinned_memory", DT_CLDEVICE_HEAD, device->config_id,
             !IS_NULL_PTR(device->cname) ? device->cname : "");
  const int pinned_memory = dt_conf_key_not_empty(key) ? dt_conf_get_int(key) : device->pinned_memory;
 
  return pinned_memory & DT_OPENCL_PINNING_ON;
}
 
int dt_opencl_set_detected_device_pinned_memory(const int detected, const gboolean enabled)
{
  const dt_opencl_detected_device_t *device = dt_opencl_get_detected_device(detected);
  if(IS_NULL_PTR(device)) return -1;
 
  gchar key[256] = { 0 };
  const int pinned_memory = enabled ? DT_OPENCL_PINNING_ON : DT_OPENCL_PINNING_OFF;
  g_snprintf(key, sizeof(key), "%s/%d/%s/pinned_memory", DT_CLDEVICE_HEAD, device->config_id,
             !IS_NULL_PTR(device->cname) ? device->cname : "");
  dt_conf_set_int(key, pinned_memory);
 
  dt_opencl_t *cl = darktable.opencl;
  cl->detected_devs[detected].pinned_memory = pinned_memory;
  // device->config_id is this device's live index into cl->dev[] (assigned at detection time,
  // see dt_opencl_device_init()) -- apply immediately instead of only on the next restart, since
  // dt_opencl_use_pinned_memory() reads cl->dev[], never cl->detected_devs[].
  if(device->config_id >= 0 && device->config_id < cl->num_devs)
    cl->dev[device->config_id].pinned_memory = pinned_memory;
  return 0;
}
 
size_t dt_opencl_detected_device_headroom(const int detected)
{
  const dt_opencl_detected_device_t *device = dt_opencl_get_detected_device(detected);
  if(IS_NULL_PTR(device)) return 0;
 
  gchar key[256] = { 0 };
  g_snprintf(key, sizeof(key), "%s/%d/%s/id%d/forced_headroom", DT_CLDEVICE_HEAD, device->config_id,
             !IS_NULL_PTR(device->cname) ? device->cname : "", device->config_id);
 
  return dt_conf_key_not_empty(key) ? (size_t)dt_conf_get_int(key) : device->forced_headroom;
}
 
int dt_opencl_set_detected_device_headroom(const int detected, const size_t headroom)
{
  const dt_opencl_detected_device_t *device = dt_opencl_get_detected_device(detected);
  if(IS_NULL_PTR(device)) return -1;
 
  gchar key[256] = { 0 };
  g_snprintf(key, sizeof(key), "%s/%d/%s/id%d/forced_headroom", DT_CLDEVICE_HEAD, device->config_id,
             !IS_NULL_PTR(device->cname) ? device->cname : "", device->config_id);
  const int clamped_headroom = (int)MIN(headroom, (size_t)G_MAXINT);
  dt_conf_set_int(key, clamped_headroom);
 
  dt_opencl_t *cl = darktable.opencl;
  cl->detected_devs[detected].forced_headroom = clamped_headroom;
  // Same live-apply as pinned memory above: device->config_id is this device's index into
  // cl->dev[], which dt_opencl_get_device_available() actually reads through used_available.
  // dt_opencl_check_tuning() recomputes that from the value we just wrote, so the new headroom
  // is in effect before this function returns instead of only after a restart.
  if(device->config_id >= 0 && device->config_id < cl->num_devs)
  {
    cl->dev[device->config_id].forced_headroom = clamped_headroom;
    dt_opencl_check_tuning(device->config_id);
  }
  return 0;
}
 
// returns 0 if all ok or an error if we failed to init this device
static int dt_opencl_device_init(dt_opencl_t *cl, const int dev, cl_device_id *devices, const int k)
{
  int res;
  cl_int err;
  gboolean lock_initialized = FALSE;
 
  memset(cl->dev[dev].program, 0x0, sizeof(cl_program) * DT_OPENCL_MAX_PROGRAMS);
  memset(cl->dev[dev].program_used, 0x0, sizeof(int) * DT_OPENCL_MAX_PROGRAMS);
  memset(cl->dev[dev].kernel, 0x0, sizeof(cl_kernel) * DT_OPENCL_MAX_KERNELS);
  memset(cl->dev[dev].kernel_used, 0x0, sizeof(int) * DT_OPENCL_MAX_KERNELS);
  cl->dev[dev].context = NULL;
  cl->dev[dev].cmd_queue = NULL;
  cl->dev[dev].eventlist = NULL;
  cl->dev[dev].eventtags = NULL;
  cl->dev[dev].numevents = 0;
  cl->dev[dev].eventsconsolidated = 0;
  cl->dev[dev].maxevents = 0;
  cl->dev[dev].maxeventslot = 0;
  cl->dev[dev].lostevents = 0;
  cl->dev[dev].totalevents = 0;
  cl->dev[dev].totalsuccess = 0;
  cl->dev[dev].totallost = 0;
  cl->dev[dev].summary = CL_COMPLETE;
  cl->dev[dev].used_global_mem = 0;
  cl->dev[dev].nvidia_sm_20 = 0;
  cl->dev[dev].vendor = NULL;
  cl->dev[dev].name = NULL;
  cl->dev[dev].cname = NULL;
  cl->dev[dev].options = NULL;
  cl->dev[dev].options_md5 = NULL;
  cl->dev[dev].memory_in_use = 0;
  cl->dev[dev].peak_memory = 0;
  cl->dev[dev].used_available = 0;
  // setting sane/conservative defaults at first
  cl->dev[dev].avoid_atomics = 0;
  cl->dev[dev].micro_nap = 250;
  cl->dev[dev].pinned_memory = DT_OPENCL_PINNING_ON;
  cl->dev[dev].clroundup_wd = 16;
  cl->dev[dev].clroundup_ht = 16;
  cl->dev[dev].use_events = 1;
  cl->dev[dev].event_handles = 128;
  cl->dev[dev].disabled = 0;
  cl->dev[dev].forced_headroom = 0;
  cl->dev[dev].runtime_error = 0;
  cl_device_id devid = cl->dev[dev].devid = devices[k];
 
  char *infostr = NULL;
  size_t infostr_size;
 
  char *cname = NULL;
  size_t cname_size;
 
  char *vendor = NULL;
  size_t vendor_size;
 
  char *driverversion = NULL;
  size_t driverversion_size;
 
  char *deviceversion = NULL;
  size_t deviceversion_size;
 
  size_t infoint;
  size_t *infointtab = NULL;
  cl_device_type type;
  cl_bool image_support = 0;
  cl_bool device_available = 0;
  cl_uint vendor_id = 0;
  cl_bool little_endian = 0;
  cl_platform_id platform_id = 0;
 
  char *dtcache = calloc(PATH_MAX, sizeof(char));
  char *cachedir = calloc(PATH_MAX, sizeof(char));
  char *devname = calloc(DT_OPENCL_CBUFFSIZE, sizeof(char));
  char *drvversion = calloc(DT_OPENCL_CBUFFSIZE, sizeof(char));
  char *platform_name = calloc(DT_OPENCL_CBUFFSIZE, sizeof(char));
  char *platform_vendor = calloc(DT_OPENCL_CBUFFSIZE, sizeof(char));
 
  char kerneldir[PATH_MAX] = { 0 };
  char *filename = calloc(PATH_MAX, sizeof(char));
  char *confentry = calloc(PATH_MAX, sizeof(char));
  char *binname = calloc(PATH_MAX, sizeof(char));
  dt_print_nts(DT_DEBUG_OPENCL, "\n[dt_opencl_device_init]\n");
 
  // test GPU availability, vendor, memory, image support etc:
  (cl->dlocl->symbols->dt_clGetDeviceInfo)(devid, CL_DEVICE_AVAILABLE, sizeof(cl_bool), &device_available, NULL);
 
  // Queried early (before dt_opencl_read_device_config() below, which needs it to pick a sane
  // default headroom on first run) rather than alongside the other capability queries further down.
  cl_bool host_unified_memory = CL_FALSE;
  (cl->dlocl->symbols->dt_clGetDeviceInfo)(devid, CL_DEVICE_HOST_UNIFIED_MEMORY, sizeof(cl_bool),
                                           &host_unified_memory, NULL);
  cl->dev[dev].host_unified_memory = (host_unified_memory == CL_TRUE);
 
  err = dt_opencl_get_device_info(cl, devid, CL_DEVICE_VENDOR, (void **)&vendor, &vendor_size);
  if(err != CL_SUCCESS)
  {
    dt_print_nts(DT_DEBUG_OPENCL, "  *** could not get vendor name of device %d: %i\n", k, err);
    res = -1;
    goto end;
  }
 
  (cl->dlocl->symbols->dt_clGetDeviceInfo)(devid, CL_DEVICE_VENDOR_ID, sizeof(cl_uint), &vendor_id, NULL);
 
  err = dt_opencl_get_device_info(cl, devid, CL_DEVICE_NAME, (void **)&infostr, &infostr_size);
  if(err != CL_SUCCESS)
  {
    dt_print_nts(DT_DEBUG_OPENCL, "  *** could not get device name of device %d: %i\n", k, err);
    res = -1;
    goto end;
  }
 
  // get the canonical device name
  cname_size = infostr_size;
  cname = malloc(cname_size);
  _ascii_str_canonical(infostr, cname, cname_size);
  cl->dev[dev].name = strdup(infostr);
  cl->dev[dev].cname = strdup(cname);
 
  // take every detected device into account of checksum
  cl->crc = crc32(cl->crc, (const unsigned char *)infostr, strlen(infostr));
 
  err = (cl->dlocl->symbols->dt_clGetDeviceInfo)(devid, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &platform_id, NULL);
  if(err != CL_SUCCESS)
  {
    g_strlcpy(platform_vendor, "no platform id", DT_OPENCL_CBUFFSIZE);
    g_strlcpy(platform_name, "no platform id", DT_OPENCL_CBUFFSIZE);
    dt_print_nts(DT_DEBUG_OPENCL, "  *** could not get platform id for device `%s' : %i\n", cl->dev[dev].name, err);
  }
  else
  {
    err = (cl->dlocl->symbols->dt_clGetPlatformInfo)(platform_id, CL_PLATFORM_NAME, DT_OPENCL_CBUFFSIZE, platform_name, NULL);
    if(err != CL_SUCCESS)
    {
      dt_print_nts(DT_DEBUG_OPENCL, "  *** could not get platform name for device `%s' : %i\n", cl->dev[dev].name, err);
      g_strlcpy(platform_name, "???", DT_OPENCL_CBUFFSIZE);
    }
 
    err = (cl->dlocl->symbols->dt_clGetPlatformInfo)(platform_id, CL_PLATFORM_VENDOR, DT_OPENCL_CBUFFSIZE, platform_vendor, NULL);
    if(err != CL_SUCCESS)
    {
      dt_print_nts(DT_DEBUG_OPENCL, "  *** could not get platform vendor for device `%s' : %i\n", cl->dev[dev].name, err);
      g_strlcpy(platform_vendor, "???", DT_OPENCL_CBUFFSIZE);
    }
  }
 
  const gboolean newdevice = dt_opencl_read_device_config(dev);
  dt_print_nts(DT_DEBUG_OPENCL, "   DEVICE:                   %d: '%s'%s\n", k, infostr, (newdevice) ? ", NEW" : "" );
  dt_print_nts(DT_DEBUG_OPENCL, "   CANONICAL NAME:           %s\n", cname);
  dt_print_nts(DT_DEBUG_OPENCL, "   PLATFORM NAME & VENDOR:   %s, %s\n", platform_name, platform_vendor);
 
  err = dt_opencl_get_device_info(cl, devid, CL_DRIVER_VERSION, (void **)&driverversion, &driverversion_size);
  if(err != CL_SUCCESS)
  {
    dt_print_nts(DT_DEBUG_OPENCL, "   *** driver version not available *** %i\n", err);
    res = -1;
    cl->dev[dev].disabled |= 1;
    goto end;
  }
 
  err = dt_opencl_get_device_info(cl, devid, CL_DEVICE_VERSION, (void **)&deviceversion, &deviceversion_size);
  if(err != CL_SUCCESS)
  {
    dt_print_nts(DT_DEBUG_OPENCL, "   *** device version not available *** %i\n", err);
    res = -1;
    cl->dev[dev].disabled |= 1;
    goto end;
  }
 
  // take every detected device driver into account of checksum
  cl->crc = crc32(cl->crc, (const unsigned char *)deviceversion, deviceversion_size);
 
  (cl->dlocl->symbols->dt_clGetDeviceInfo)(devid, CL_DEVICE_TYPE, sizeof(cl_device_type), &type, NULL);
  (cl->dlocl->symbols->dt_clGetDeviceInfo)(devid, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool), &image_support, NULL);
  (cl->dlocl->symbols->dt_clGetDeviceInfo)(devid, CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof(size_t),
                                           &(cl->dev[dev].max_image_height), NULL);
  (cl->dlocl->symbols->dt_clGetDeviceInfo)(devid, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof(size_t),
                                           &(cl->dev[dev].max_image_width), NULL);
  (cl->dlocl->symbols->dt_clGetDeviceInfo)(devid, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(cl_ulong),
                                           &(cl->dev[dev].max_mem_alloc), NULL);
  (cl->dlocl->symbols->dt_clGetDeviceInfo)(devid, CL_DEVICE_ENDIAN_LITTLE, sizeof(cl_bool), &little_endian, NULL);
 
  cl->dev[dev].cltype = (unsigned int)type;
 
 
  if(!strncasecmp(vendor, "NVIDIA", 6))
  {
    // very lame attempt to detect support for atomic float add in global memory.
    // we need compute model sm_20, but let's try for all nvidia devices :(
    cl->dev[dev].nvidia_sm_20 = dt_nvidia_gpu_supports_sm_20(infostr);
  }
 
  const gboolean is_cpu_device = (type & CL_DEVICE_TYPE_CPU) == CL_DEVICE_TYPE_CPU;
 
  // micro_nap can be made less conservative on current systems at least if not on-CPU
  if(newdevice)
    cl->dev[dev].micro_nap = (is_cpu_device) ? 1000 : 250;
 
  dt_print_nts(DT_DEBUG_OPENCL, "   DRIVER VERSION:           %s\n", driverversion);
  dt_print_nts(DT_DEBUG_OPENCL, "   DEVICE VERSION:           %s%s\n", deviceversion,
     cl->dev[dev].nvidia_sm_20 ? ", SM_20 SUPPORT" : "");
  dt_print_nts(DT_DEBUG_OPENCL, "   DEVICE_TYPE:              %s%s%s\n",
      ((type & CL_DEVICE_TYPE_CPU) == CL_DEVICE_TYPE_CPU) ? "CPU" : "",
      ((type & CL_DEVICE_TYPE_GPU) == CL_DEVICE_TYPE_GPU) ? "GPU" : "",
      (type & CL_DEVICE_TYPE_ACCELERATOR)                 ? ", Accelerator" : "" );
 
  if(is_cpu_device && newdevice)
  {
    dt_print_nts(DT_DEBUG_OPENCL, "   *** discarding new device as emulated by CPU ***\n");
    cl->dev[dev].disabled |= 1;
    res = -1;
    goto end;
  }
 
  if(!device_available)
  {
    dt_print_nts(DT_DEBUG_OPENCL, "   *** device is not available ***\n");
    res = -1;
    goto end;
  }
 
  if(!image_support)
  {
    dt_print_nts(DT_DEBUG_OPENCL, "   *** The OpenCL driver doesn't provide image support. See also 'clinfo' output ***\n");
    res = -1;
    cl->dev[dev].disabled |= 1;
    goto end;
  }
 
  if(!little_endian)
  {
    dt_print_nts(DT_DEBUG_OPENCL, "   *** device is not little endian ***\n");
    res = -1;
    cl->dev[dev].disabled |= 1;
    goto end;
  }
 
  (cl->dlocl->symbols->dt_clGetDeviceInfo)(devid, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(cl_ulong),
                                           &(cl->dev[dev].max_global_mem), NULL);
  if(cl->dev[dev].max_global_mem < (uint64_t)512ul * 1024ul * 1024ul)
  {
    dt_print_nts(DT_DEBUG_OPENCL, "   *** insufficient global memory (%" PRIu64 "MB) ***\n",
                                   cl->dev[dev].max_global_mem / 1024 / 1024);
    res = -1;
    cl->dev[dev].disabled |= 1;
    goto end;
  }
 
  cl->dev[dev].vendor = strdup(dt_opencl_get_vendor_by_id(vendor_id));
 
  const gboolean is_blacklisted = dt_opencl_check_driver_blacklist(deviceversion);
 
  // disable device for now if this is the first time detected and blacklisted too.
  if(newdevice && is_blacklisted)
  {
    // To keep installations we look for the old blacklist conf key
    const gboolean old_blacklist = dt_conf_get_bool("opencl_disable_drivers_blacklist");
    cl->dev[dev].disabled |= (old_blacklist) ? 0 : 1;
    if(cl->dev[dev].disabled)
      dt_print_nts(DT_DEBUG_OPENCL, "   *** new device is blacklisted ***\n");
    res = -1;
    goto end;
  }
 
  dt_print_nts(DT_DEBUG_OPENCL, "   GLOBAL MEM SIZE:          %.0f MB\n", (double)cl->dev[dev].max_global_mem / 1024.0 / 1024.0);
  dt_print_nts(DT_DEBUG_OPENCL, "   MAX MEM ALLOC:            %.0f MB\n", (double)cl->dev[dev].max_mem_alloc / 1024.0 / 1024.0);
  dt_print_nts(DT_DEBUG_OPENCL, "   MAX IMAGE SIZE:           %" G_GSIZE_FORMAT " x %" G_GSIZE_FORMAT "\n", cl->dev[dev].max_image_width, cl->dev[dev].max_image_height);
  (cl->dlocl->symbols->dt_clGetDeviceInfo)(devid, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(infoint), &infoint, NULL);
  dt_print_nts(DT_DEBUG_OPENCL, "   MAX WORK GROUP SIZE:      %" G_GSIZE_FORMAT "\n", infoint);
  (cl->dlocl->symbols->dt_clGetDeviceInfo)(devid, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, sizeof(infoint), &infoint, NULL);
  dt_print_nts(DT_DEBUG_OPENCL, "   MAX WORK ITEM DIMENSIONS: %" G_GSIZE_FORMAT "\n", infoint);
 
  size_t infointtab_size;
  err = dt_opencl_get_device_info(cl, devid, CL_DEVICE_MAX_WORK_ITEM_SIZES, (void **)&infointtab, &infointtab_size);
  if(err == CL_SUCCESS)
  {
    dt_print_nts(DT_DEBUG_OPENCL, "   MAX WORK ITEM SIZES:      [ ");
    for(size_t i = 0; i < infoint; i++) dt_print_nts(DT_DEBUG_OPENCL, "%" G_GSIZE_FORMAT " ", infointtab[i]);
    dt_free(infointtab);
    dt_print_nts(DT_DEBUG_OPENCL, "]\n");
  }
  else
  {
    dt_print_nts(DT_DEBUG_OPENCL, "   *** could not get maximum work item sizes ***\n");
    res = -1;
    cl->dev[dev].disabled |= 1;
    goto end;
  }
 
  const gboolean pinning = (cl->dev[dev].pinned_memory & DT_OPENCL_PINNING_ON);
  dt_print_nts(DT_DEBUG_OPENCL, "   PINNED MEMORY TRANSFER:   %s\n", pinning ? "WANTED" : "NO");
  dt_print_nts(DT_DEBUG_OPENCL, "   FORCED HEADROOM:          %" G_GSIZE_FORMAT "\n", cl->dev[dev].forced_headroom);
  dt_print_nts(DT_DEBUG_OPENCL, "   AVOID ATOMICS:            %s\n", (cl->dev[dev].avoid_atomics) ? "YES" : "NO");
  dt_print_nts(DT_DEBUG_OPENCL, "   MICRO NAP:                %i\n", cl->dev[dev].micro_nap);
  dt_print_nts(DT_DEBUG_OPENCL, "   ROUNDUP WIDTH:            %i\n", cl->dev[dev].clroundup_wd);
  dt_print_nts(DT_DEBUG_OPENCL, "   ROUNDUP HEIGHT:           %i\n", cl->dev[dev].clroundup_ht);
  dt_print_nts(DT_DEBUG_OPENCL, "   CHECK EVENT HANDLES:      %i\n", cl->dev[dev].event_handles);
  dt_print_nts(DT_DEBUG_OPENCL, "   DEFAULT DEVICE:           %s\n", (type & CL_DEVICE_TYPE_DEFAULT) ? "YES" : "NO");
 
  if(type & CL_DEVICE_TYPE_GPU)
  {
    dt_opencl_detected_device_t *detected_devs
        = g_realloc(cl->detected_devs, sizeof(*cl->detected_devs) * (cl->num_detected_devs + 1));
    if(!IS_NULL_PTR(detected_devs))
    {
      cl->detected_devs = detected_devs;
      dt_opencl_detected_device_t *detected = cl->detected_devs + cl->num_detected_devs;
      detected->config_id = dev;
      detected->name = g_strdup(cl->dev[dev].name);
      detected->cname = g_strdup(cl->dev[dev].cname);
      detected->cltype = cl->dev[dev].cltype;
      detected->disabled = cl->dev[dev].disabled & 1;
      detected->pinned_memory = cl->dev[dev].pinned_memory;
      detected->forced_headroom = cl->dev[dev].forced_headroom;
      detected->host_unified_memory = cl->dev[dev].host_unified_memory;
      cl->num_detected_devs++;
    }
  }
 
  if(cl->dev[dev].disabled)
  {
    dt_print_nts(DT_DEBUG_OPENCL, "   *** marked as disabled ***\n");
    res = -1;
    goto end;
  }
  dt_print_nts(DT_DEBUG_OPENCL, "   *** Device enabled ***\n");
 
  dt_pthread_mutex_init(&cl->dev[dev].lock, NULL);
  lock_initialized = TRUE;
 
  cl->dev[dev].context = (cl->dlocl->symbols->dt_clCreateContext)(0, 1, &devid, NULL, NULL, &err);
  if(err != CL_SUCCESS)
  {
    dt_print_nts(DT_DEBUG_OPENCL, "   *** could not create context *** %i\n", err);
    res = -1;
    goto end;
  }
  // create a command queue for first device the context reported
  cl->dev[dev].cmd_queue = (cl->dlocl->symbols->dt_clCreateCommandQueue)(
      cl->dev[dev].context, devid, (darktable.unmuted & DT_DEBUG_PERF) ? CL_QUEUE_PROFILING_ENABLE : 0, &err);
  if(err != CL_SUCCESS)
  {
    dt_print_nts(DT_DEBUG_OPENCL, "   *** could not create command queue *** %i\n", err);
    res = -1;
    goto end;
  }
 
  dt_loc_get_kerneldir(kerneldir, sizeof(kerneldir));
  dt_print_nts(DT_DEBUG_OPENCL, "   KERNEL SOURCE DIRECTORY:  %s\n", kerneldir);
 
  double tstart, tend, tdiff;
  dt_loc_get_user_cache_dir(dtcache, PATH_MAX * sizeof(char));
 
  int len = MIN(strlen(infostr),1024 * sizeof(char));;
  int j = 0;
  // remove non-alphanumeric chars from device name
  for(int i = 0; i < len; i++)
    if(isalnum(infostr[i])) devname[j++] = infostr[i];
  devname[j] = 0;
  len = MIN(strlen(driverversion), 1024 * sizeof(char));
  j = 0;
  // remove non-alphanumeric chars from driver version
  for(int i = 0; i < len; i++)
    if(isalnum(driverversion[i])) drvversion[j++] = driverversion[i];
  drvversion[j] = 0;
  snprintf(cachedir, PATH_MAX * sizeof(char), "%s" G_DIR_SEPARATOR_S "cached_kernels_for_%s_%s", dtcache, devname, drvversion);
 
  dt_print_nts(DT_DEBUG_OPENCL, "   KERNEL BUILD DIRECTORY:   %s\n", cachedir);
 
  if(g_mkdir_with_parents(cachedir, 0700) == -1)
  {
    dt_print_nts(DT_DEBUG_OPENCL, "   *** failed to create kernel directory `%s' ***\n", cachedir);
    res = -1;
    goto end;
  }
 
  dt_concat_path_file(filename, kerneldir, "programs.conf");
 
  char *escapedkerneldir = NULL;
#ifndef __APPLE__
  escapedkerneldir = g_strdup_printf("\"%s\"", kerneldir);
#else
  escapedkerneldir = dt_util_str_replace(kerneldir, " ", "\\ ");
#endif
 
  gchar* compile_option_name_cname = g_strdup_printf("%s/%i/%s/building", DT_CLDEVICE_HEAD, dev, cl->dev[dev].cname);
  const char* compile_opt = NULL;
 
  if(dt_conf_key_exists(compile_option_name_cname))
    compile_opt = dt_conf_get_string_const(compile_option_name_cname);
  else
  {
    switch(vendor_id)
    {
      case DT_OPENCL_VENDOR_AMD:
        compile_opt = DT_OPENCL_DEFAULT_COMPILE_AMD;
        break;
      case DT_OPENCL_VENDOR_NVIDIA:
        compile_opt = DT_OPENCL_DEFAULT_COMPILE_NVIDIA;
        break;
      case DT_OPENCL_VENDOR_INTEL:
        compile_opt = DT_OPENCL_DEFAULT_COMPILE_INTEL;
        break;
      default:
        compile_opt = DT_OPENCL_DEFAULT_COMPILE;
    }
  }
  gchar *my_option = g_strdup(compile_opt);
  dt_conf_set_string(compile_option_name_cname, my_option);
 
  cl->dev[dev].options = g_strdup_printf("-w %s %s -D%s=1 -I%s",
                            my_option,
                            (cl->dev[dev].nvidia_sm_20 ? " -DNVIDIA_SM_20=1" : ""),
                            dt_opencl_get_vendor_by_id(vendor_id), escapedkerneldir);
  // Keep kernel checksum stable when the runtime kernel path changes (e.g. AppImage mount point).
  const char *kerneldir_token = "<ansel-kernels>";
  char *escapedkerneldir_md5 = NULL;
#ifndef __APPLE__
  escapedkerneldir_md5 = g_strdup_printf("\"%s\"", kerneldir_token);
#else
  escapedkerneldir_md5 = g_strdup(kerneldir_token);
#endif
  cl->dev[dev].options_md5 = g_strdup_printf("-w %s %s -D%s=1 -I%s",
                               my_option,
                               (cl->dev[dev].nvidia_sm_20 ? " -DNVIDIA_SM_20=1" : ""),
                               dt_opencl_get_vendor_by_id(vendor_id), escapedkerneldir_md5);
 
  dt_print_nts(DT_DEBUG_OPENCL, "   CL COMPILER OPTION:       %s\n", my_option);
 
  dt_free(compile_option_name_cname);
  dt_free(my_option);
  dt_free(escapedkerneldir);
  dt_free(escapedkerneldir_md5);
  escapedkerneldir = NULL;
 
  const char *clincludes[DT_OPENCL_MAX_INCLUDES] = { "rgb_norms.h", "noise_generator.h", "color_conversion.h", "colorspaces.cl", "colorspace.h", "common.h", NULL };
  char *includemd5[DT_OPENCL_MAX_INCLUDES] = { NULL };
  dt_opencl_md5sum(clincludes, includemd5);
 
  if(newdevice) // so far the device seems to be ok. Make sure to write&export the conf database to
  {
    dt_opencl_write_device_config(dev);
    dt_conf_save(darktable.conf);
  }
 
  // now load all darktable cl kernels.
  // TODO: compile as a job?
  tstart = dt_get_wtime();
  FILE *f = g_fopen(filename, "rb");
  if(f)
  {
    while(!feof(f))
    {
      int prog = -1;
      gchar *confline_pattern = g_strdup_printf("%%%" G_GSIZE_FORMAT "[^\n]\n", PATH_MAX * sizeof(char) - 1);
      int rd = fscanf(f, confline_pattern, confentry);
      dt_free(confline_pattern);
      if(rd != 1) continue;
      // remove comments:
      size_t end = strlen(confentry);
      for(size_t pos = 0; pos < end; pos++)
        if(confentry[pos] == '#')
        {
          confentry[pos] = '\0';
          for(int l = pos - 1; l >= 0; l--)
          {
            if(confentry[l] == ' ')
              confentry[l] = '\0';
            else
              break;
          }
          break;
        }
      if(confentry[0] == '\0') continue;
 
      const char *programname = NULL, *programnumber = NULL;
      gchar **tokens = g_strsplit_set(confentry, " \t", 2);
      if(tokens)
      {
        programname = tokens[0];
        if(tokens[0])
          programnumber = tokens[1]; // if the 0st wasn't NULL then we have at least the terminating NULL in [1]
      }
 
      prog = programnumber ? strtol(programnumber, NULL, 10) : -1;
 
      if(IS_NULL_PTR(programname) || programname[0] == '\0' || prog < 0)
      {
        dt_print(DT_DEBUG_OPENCL, "[dt_opencl_device_init] malformed entry in programs.conf `%s'; ignoring it!\n", confentry);
        continue;
      }
      dt_concat_path_file(filename, kerneldir, programname);
      gchar *program_bin = g_strdup_printf("%s.bin", programname);
      dt_concat_path_file(binname, cachedir, program_bin);
      dt_free(program_bin);
 
      dt_vprint(DT_DEBUG_OPENCL, "[dt_opencl_device_init] testing program `%s' ..\n", programname);
      int loaded_cached;
      char md5sum[33];
      if(dt_opencl_load_program(dev, prog, filename, binname, cachedir, md5sum, includemd5, &loaded_cached))
      {
        if(!loaded_cached)
          _opencl_splash_update_compile(programname);
 
        if(dt_opencl_build_program(dev, prog, binname, cachedir, md5sum, loaded_cached) != CL_SUCCESS)
        {
          dt_print(DT_DEBUG_OPENCL, "[dt_opencl_device_init] failed to compile program `%s'!\n", programname);
          fclose(f);
          g_strfreev(tokens);
          res = -1;
          goto end;
        }
      }
 
      g_strfreev(tokens);
    }
 
    fclose(f);
    tend = dt_get_wtime();
    tdiff = tend - tstart;
    dt_print_nts(DT_DEBUG_OPENCL, "   KERNEL LOADING TIME:       %2.4lf sec\n", tdiff);
  }
  else
  {
    dt_print_nts(DT_DEBUG_OPENCL, "[dt_opencl_device_init] could not open `%s'!\n", filename);
    res = -1;
    goto end;
  }
  for(int n = 0; n < DT_OPENCL_MAX_INCLUDES; n++) dt_free(includemd5[n]);
  res = 0;
 
end:
  // we always write the device config to keep track of disabled devices
  dt_opencl_write_device_config(dev);
 
  if(res != 0)
  {
    if(lock_initialized)
    {
      for(int n = 0; n < DT_OPENCL_MAX_KERNELS; n++)
        if(cl->dev[dev].kernel_used[n]) (cl->dlocl->symbols->dt_clReleaseKernel)(cl->dev[dev].kernel[n]);
      for(int n = 0; n < DT_OPENCL_MAX_PROGRAMS; n++)
        if(cl->dev[dev].program_used[n]) (cl->dlocl->symbols->dt_clReleaseProgram)(cl->dev[dev].program[n]);
      if(!IS_NULL_PTR(cl->dev[dev].cmd_queue))
        (cl->dlocl->symbols->dt_clReleaseCommandQueue)(cl->dev[dev].cmd_queue);
      if(!IS_NULL_PTR(cl->dev[dev].context))
        (cl->dlocl->symbols->dt_clReleaseContext)(cl->dev[dev].context);
      dt_pthread_mutex_destroy(&cl->dev[dev].lock);
    }
 
    dt_free(cl->dev[dev].vendor);
    dt_free(cl->dev[dev].name);
    dt_free(cl->dev[dev].cname);
    dt_free(cl->dev[dev].options);
    dt_free(cl->dev[dev].options_md5);
  }
 
  dt_free(infostr);
  dt_free(cname);
  dt_free(vendor);
  dt_free(driverversion);
  dt_free(deviceversion);
 
  dt_free(dtcache);
  dt_free(cachedir);
  dt_free(devname);
  dt_free(drvversion);
  dt_free(platform_name);
  dt_free(platform_vendor);
 
  dt_free(filename);
  dt_free(confentry);
  dt_free(binname);
 
  return res;
}
 
void dt_opencl_init(dt_opencl_t *cl, const gboolean exclude_opencl, const gboolean print_statistics)
{
  _opencl_splash_active = FALSE;
 
  dt_pthread_mutex_init(&cl->lock, NULL);
  dt_pthread_mutex_init(&cl->mem_sizes_lock, NULL);
  cl->mem_sizes = g_hash_table_new_full(g_direct_hash, g_direct_equal, NULL, g_free);
  cl->inited = 0;
  cl->enabled = 0;
  cl->stopped = 0;
  cl->error_count = 0;
  cl->print_statistics = print_statistics;
 
  // work-around to fix a bug in some AMD OpenCL compilers, which would fail parsing certain numerical
  // constants if locale is different from "C".
  // we save the current locale, set locale to "C", and restore the previous setting after OpenCL is
  // initialized
  char *locale = strdup(setlocale(LC_ALL, NULL));
  setlocale(LC_ALL, "C");
 
  cl->crc = 5781;
  cl->dlocl = NULL;
  cl->dev_priority_image = 0;
  cl->dev_priority_preview = 0;
  cl->dev_priority_export = 0;
  cl->dev_priority_thumbnail = 0;
  cl->num_detected_devs = 0;
  cl->detected_devs = NULL;
 
  if(exclude_opencl) return;
 
  cl_platform_id *all_platforms = NULL;
  cl_uint *all_num_devices = NULL;
 
  char *platform_name = calloc(DT_OPENCL_CBUFFSIZE, sizeof(char));
  char *platform_vendor = calloc(DT_OPENCL_CBUFFSIZE, sizeof(char));
 
  dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] opencl related configuration options:\n");
  dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] opencl: %s\n", dt_conf_get_bool("opencl") ? "ON" : "OFF" );
  // look for explicit definition of opencl_runtime library in preferences
  const char *library = dt_conf_get_string_const("opencl_library");
  dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] opencl_library: '%s'\n", (strlen(library) == 0) ? "default path" : library);
  dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] opencl_mandatory_timeout: %d\n",
           dt_conf_get_int("opencl_mandatory_timeout"));
 
  // dynamically load opencl runtime
  if((cl->dlocl = dt_dlopencl_init(library)) == NULL)
  {
    dt_print_nts(DT_DEBUG_OPENCL,
             "[opencl_init] no working opencl library found. Continue with opencl disabled\n");
    goto finally;
  }
  else
  {
    dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] opencl library '%s' found on your system and loaded\n",
             cl->dlocl->library);
  }
 
  cl_int err;
  all_platforms = malloc(sizeof(cl_platform_id) * DT_OPENCL_MAX_PLATFORMS);
  all_num_devices = malloc(sizeof(cl_uint) * DT_OPENCL_MAX_PLATFORMS);
  cl_uint num_platforms = DT_OPENCL_MAX_PLATFORMS;
  err = (cl->dlocl->symbols->dt_clGetPlatformIDs)(DT_OPENCL_MAX_PLATFORMS, all_platforms, &num_platforms);
  if(err != CL_SUCCESS)
  {
    dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] could not get platforms: %i\n", err);
    goto finally;
  }
 
  if(num_platforms == 0)
  {
    dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] no opencl platform available\n");
    goto finally;
  }
  dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] found %d platform%s\n", num_platforms,
           num_platforms > 1 ? "s" : "");
 
  for(int n = 0; n < num_platforms; n++)
  {
    cl_platform_id platform = all_platforms[n];
    // get the number of GPU devices available to the platforms
    // the other common option is CL_DEVICE_TYPE_GPU/CPU (but the latter doesn't work with the nvidia drivers)
    err = (cl->dlocl->symbols->dt_clGetDeviceIDs)(platform, CL_DEVICE_TYPE_ALL, 0, NULL, &(all_num_devices[n]));
    if(err != CL_SUCCESS)
    {
      cl_int errv = (cl->dlocl->symbols->dt_clGetPlatformInfo)(platform, CL_PLATFORM_VENDOR, DT_OPENCL_CBUFFSIZE, platform_vendor, NULL);
      cl_int errn = (cl->dlocl->symbols->dt_clGetPlatformInfo)(platform, CL_PLATFORM_NAME, DT_OPENCL_CBUFFSIZE, platform_name, NULL);
      if((errn == CL_SUCCESS) && (errv == CL_SUCCESS))
        dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] no devices found for %s (vendor) - %s (name)\n", platform_vendor, platform_name);
      else
        dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] no devices found for unknown platform\n");
 
      all_num_devices[n] = 0;
    }
    else
    {
      char profile[64] = { 0 };
      size_t profile_size;
      err = (cl->dlocl->symbols->dt_clGetPlatformInfo)(platform, CL_PLATFORM_PROFILE, 64, profile, &profile_size);
      if(err != CL_SUCCESS)
      {
        all_num_devices[n] = 0;
        dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] could not get profile: %i\n", err);
      }
      else
      {
        // fprintf(stderr, "%s\n", profile);
        if(strcmp("FULL_PROFILE", profile) != 0)
        {
          all_num_devices[n] = 0;
          dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] platform %i is not FULL_PROFILE\n", n);
        }
      }
    }
  }
 
  cl_uint num_devices = 0;
  for(int n = 0; n < num_platforms; n++) num_devices += all_num_devices[n];
 
  // create the device list
  cl_device_id *devices = 0;
  if(num_devices)
  {
    cl->dev = (dt_opencl_device_t *)malloc(sizeof(dt_opencl_device_t) * num_devices);
    devices = (cl_device_id *)malloc(sizeof(cl_device_id) * num_devices);
    if(IS_NULL_PTR(cl->dev) || IS_NULL_PTR(devices))
    {
      dt_free(cl->dev);
      dt_free(devices);
      dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] could not allocate memory\n");
      goto finally;
    }
  }
 
  cl_device_id *devs = devices;
  for(int n = 0; n < num_platforms; n++)
  {
    if(all_num_devices[n])
    {
      cl_platform_id platform = all_platforms[n];
      err = (cl->dlocl->symbols->dt_clGetDeviceIDs)(platform, CL_DEVICE_TYPE_ALL, all_num_devices[n], devs,
                                                    NULL);
      if(err != CL_SUCCESS)
      {
        num_devices -= all_num_devices[n];
        dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] could not get devices list: %i\n", err);
      }
      devs += all_num_devices[n];
    }
  }
  devs = NULL;
 
  dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] found %d device%s\n", num_devices, num_devices > 1 ? "s" : "");
  if(num_devices == 0)
  {
    if(devices)
    {
      dt_free(devices);
    }
    goto finally;
  }
 
  int dev = 0;
  for(int k = 0; k < num_devices; k++)
  {
    const int res = dt_opencl_device_init(cl, dev, devices, k);
    if(res != 0)
      continue;
    // increase dev only if dt_opencl_device_init was successful (res == 0)
    ++dev;
  }
  dt_free(devices);
 
  if(dev > 0)
  {
    cl->num_devs = dev;
    cl->inited = 1;
    cl->enabled = dt_conf_get_bool("opencl");
    memset(cl->mandatory, 0, sizeof(cl->mandatory));
    cl->dev_priority_image = (int *)malloc(sizeof(int) * (dev + 1));
    cl->dev_priority_preview = (int *)malloc(sizeof(int) * (dev + 1));
    cl->dev_priority_export = (int *)malloc(sizeof(int) * (dev + 1));
    cl->dev_priority_thumbnail = (int *)malloc(sizeof(int) * (dev + 1));
 
    // only check successful malloc in debug mode; darktable will crash anyhow sooner or later if mallocs that
    // small would fail
    assert(!IS_NULL_PTR(cl->dev_priority_image) && !IS_NULL_PTR(cl->dev_priority_preview)
           && !IS_NULL_PTR(cl->dev_priority_export) && !IS_NULL_PTR(cl->dev_priority_thumbnail));
 
    dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] OpenCL successfully initialized.\n");
    dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] here are the internal numbers and names of OpenCL devices available to Ansel:\n");
    for(int i = 0; i < dev; i++) dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init]\t\t%d\t'%s'\n", i, cl->dev[i].name);
  }
  else
  {
    dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] no suitable devices found.\n");
  }
 
finally:
  dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] FINALLY: opencl is %sAVAILABLE on this system.\n",
           cl->inited ? "" : "NOT ");
  dt_print_nts(DT_DEBUG_OPENCL, "[opencl_init] initial status of opencl enabled flag is %s.\n",
           cl->enabled ? "ON" : "OFF");
 
  char checksum[64];
  snprintf(checksum, sizeof(checksum), "%u", cl->crc);
 
  if(cl->inited)
  {
    dt_capabilities_add("opencl");
    cl->blendop = dt_develop_blend_init_cl_global();
    cl->bilateral = dt_bilateral_init_cl_global();
    cl->gaussian = dt_gaussian_init_cl_global();
    cl->interpolation = dt_interpolation_init_cl_global();
    cl->local_laplacian = dt_local_laplacian_init_cl_global();
    cl->dwt = dt_dwt_init_cl_global();
    cl->heal = dt_heal_init_cl_global();
    cl->colorspaces = dt_colorspaces_init_cl_global();
    cl->guided_filter = dt_guided_filter_init_cl_global();
  }
 
  dt_opencl_apply_scheduling_profile();
 
  if(!cl->inited)// initialization failed
  {
    for(int i = 0; cl->dev && i < cl->num_devs; i++) dt_opencl_cleanup_device(cl, i);
  }
 
  dt_free(all_num_devices);
  dt_free(all_platforms);
  dt_free(platform_name);
  dt_free(platform_vendor);
 
  if(locale)
  {
    setlocale(LC_ALL, locale);
    dt_free(locale);
  }
 
  return;
}
 
void dt_opencl_cleanup_device(dt_opencl_t *cl, int i)
{
  dt_pthread_mutex_destroy(&cl->dev[i].lock);
  for(int k = 0; k < DT_OPENCL_MAX_KERNELS; k++)
    if(cl->dev[i].kernel_used[k]) (cl->dlocl->symbols->dt_clReleaseKernel)(cl->dev[i].kernel[k]);
  for(int k = 0; k < DT_OPENCL_MAX_PROGRAMS; k++)
    if(cl->dev[i].program_used[k]) (cl->dlocl->symbols->dt_clReleaseProgram)(cl->dev[i].program[k]);
  if(!IS_NULL_PTR(cl->dev[i].cmd_queue))
    (cl->dlocl->symbols->dt_clReleaseCommandQueue)(cl->dev[i].cmd_queue);
  if(!IS_NULL_PTR(cl->dev[i].context))
    (cl->dlocl->symbols->dt_clReleaseContext)(cl->dev[i].context);
 
  if(cl->print_statistics && (darktable.unmuted & DT_DEBUG_MEMORY))
  {
    dt_print_nts(DT_DEBUG_OPENCL, " [opencl_summary_statistics] device '%s' (%d): peak memory usage %" G_GSIZE_FORMAT " bytes (%.1f MB)\n",
                cl->dev[i].name, i, cl->dev[i].peak_memory, (float)cl->dev[i].peak_memory/(1024*1024));
  }
 
  if(cl->print_statistics && cl->dev[i].use_events)
  {
    if(cl->dev[i].totalevents)
    {
      dt_print_nts(DT_DEBUG_OPENCL, " [opencl_summary_statistics] device '%s' (%d): %d out of %d events were "
                                "successful and %d events lost. max event=%d%s\n",
        cl->dev[i].name, i, cl->dev[i].totalsuccess, cl->dev[i].totalevents, cl->dev[i].totallost,
        cl->dev[i].maxeventslot, (cl->dev[i].maxeventslot > 1024) ? "\n *** Warning, slots > 1024" : "");
    }
    else
    {
      dt_print_nts(DT_DEBUG_OPENCL, " [opencl_summary_statistics] device '%s' (%d): NOT utilized\n",
                cl->dev[i].name, i);
    }
  }
 
  if(cl->dev[i].use_events)
  {
    dt_opencl_events_reset(i);
 
    dt_free(cl->dev[i].eventlist);
    dt_free(cl->dev[i].eventtags);
  }
 
  dt_free(cl->dev[i].vendor);
  dt_free(cl->dev[i].name);
  dt_free(cl->dev[i].cname);
  dt_free(cl->dev[i].options);
  dt_free(cl->dev[i].options_md5);
}
 
void dt_opencl_cleanup(dt_opencl_t *cl)
{
  if(cl->inited)
  {
    dt_develop_blend_free_cl_global(cl->blendop);
    dt_bilateral_free_cl_global(cl->bilateral);
    dt_gaussian_free_cl_global(cl->gaussian);
    dt_interpolation_free_cl_global(cl->interpolation);
    dt_local_laplacian_free_cl_global(cl->local_laplacian);
    dt_dwt_free_cl_global(cl->dwt);
    dt_heal_free_cl_global(cl->heal);
    dt_colorspaces_free_cl_global(cl->colorspaces);
    dt_guided_filter_free_cl_global(cl->guided_filter);
 
    for(int i = 0; i < cl->num_devs; i++)
      dt_opencl_cleanup_device(cl, i);
 
    dt_free(cl->dev_priority_image);
    dt_free(cl->dev_priority_preview);
    dt_free(cl->dev_priority_export);
    dt_free(cl->dev_priority_thumbnail);
  }
 
  if(cl->dlocl)
  {
    dt_free(cl->dlocl->symbols);
    dt_free(cl->dlocl->library);
    dt_free(cl->dlocl);
  }
 
  for(int i = 0; i < cl->num_detected_devs; i++)
  {
    dt_free(cl->detected_devs[i].name);
    dt_free(cl->detected_devs[i].cname);
  }
  dt_free(cl->detected_devs);
 
  dt_free(cl->dev);
  if(cl->mem_sizes) g_hash_table_destroy(cl->mem_sizes);
  dt_pthread_mutex_destroy(&cl->mem_sizes_lock);
  dt_pthread_mutex_destroy(&cl->lock);
}
 
static const char *dt_opencl_get_vendor_by_id(unsigned int id)
{
  const char *vendor;
 
  switch(id)
  {
    case DT_OPENCL_VENDOR_AMD:
      vendor = "AMD";
      break;
    case DT_OPENCL_VENDOR_NVIDIA:
      vendor = "NVIDIA";
      break;
    case DT_OPENCL_VENDOR_INTEL:
      vendor = "INTEL";
      break;
    default:
      vendor = "UNKNOWN";
  }
 
  return vendor;
}
 
gboolean dt_opencl_finish(const int devid)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited || devid < 0) return FALSE;
 
  cl_int err = (cl->dlocl->symbols->dt_clFinish)(cl->dev[devid].cmd_queue);
 
  // take the opportunity to release some event handles, but without printing
  // summary statistics
  cl_int success = dt_opencl_events_flush(devid, 0);
 
  return (err == CL_SUCCESS && success == CL_COMPLETE);
}
 
int dt_opencl_enqueue_barrier(const int devid)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited || devid < 0) return -1;
  return (cl->dlocl->symbols->dt_clEnqueueBarrier)(cl->dev[devid].cmd_queue);
}
 
static int _take_from_list(int *list, int value)
{
  int result = -1;
 
  while(*list != -1 && *list != value) list++;
  result = *list;
 
  while(*list != -1)
  {
    *list = *(list + 1);
    list++;
  }
 
  return result;
}
 
 
static int _device_by_cname(const char *name)
{
  dt_opencl_t *cl = darktable.opencl;
  int devs = cl->num_devs;
  char tmp[2048] = { 0 };
  int result = -1;
 
  _ascii_str_canonical(name, tmp, sizeof(tmp));
 
  for(int i = 0; i < devs; i++)
  {
    if(!strcmp(tmp, cl->dev[i].cname))
    {
      result = i;
      break;
    }
  }
 
  return result;
}
 
 
static char *_ascii_str_canonical(const char *in, char *out, int maxlen)
{
  if(IS_NULL_PTR(out))
  {
    maxlen = strlen(in) + 1;
    out = malloc(maxlen);
    if(IS_NULL_PTR(out)) return NULL;
  }
 
  int len = 0;
 
  while(*in != '\0' && len < maxlen - 1)
  {
    int n = strcspn(in, "0123456789abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ");
    in += n;
    if(n != 0) continue;
    out[len] = tolower(*in);
    len++;
    in++;
  }
  out[len] = '\0';
 
  return out;
}
 
// parse a single token of priority string and store priorities in priority_list
static void dt_opencl_priority_parse(dt_opencl_t *cl, char *configstr, int *priority_list, int *mandatory)
{
  int devs = cl->num_devs;
  int count = 0;
  int *full = malloc(sizeof(int) * (devs + 1));
  int mnd = 0;
 
  // NULL or empty configstring?
  if(IS_NULL_PTR(configstr) || *configstr == '\0')
  {
    priority_list[0] = -1;
    *mandatory = 0;
    dt_free(full);
    return;
  }
 
  // check if user wants us to force-use opencl device(s)
  if(configstr[0] == '+')
  {
    mnd = 1;
    configstr++;
  }
 
  // first start with a full list of devices to take from
  for(int i = 0; i < devs; i++) full[i] = i;
  full[devs] = -1;
 
  gchar **tokens = g_strsplit(configstr, ",", 0);
  gchar **tokens_ptr = tokens;
 
  while(!IS_NULL_PTR(tokens) && !IS_NULL_PTR(*tokens_ptr) && count < devs + 1 && full[0] != -1)
  {
    gchar *str = *tokens_ptr;
    int not = 0;
    int all = 0;
 
    switch(*str)
    {
      case '*':
        all = 1;
        break;
      case '!':
        not = 1;
        while(*str == '!') str++;
        break;
    }
 
    if(all)
    {
      // copy all remaining device numbers from full to priority list
      for(int i = 0; i < devs && full[i] != -1; i++)
      {
        priority_list[count] = full[i];
        count++;
      }
      full[0] = -1; // mark full list as empty
    }
    else if(*str != '\0')
    {
      char *endptr = NULL;
 
      // first check if str corresponds to an existing canonical device name
      long number = _device_by_cname(str);
 
      // if not try to convert string into decimal device number
      if(number < 0) number = strtol(str, &endptr, 10);
 
      // still not found or negative number given? set number to -1
      if(number < 0 || (number == 0 && endptr == str)) number = -1;
 
      // try to take number out of remaining device list
      int dev_number = _take_from_list(full, number);
 
      if(!not&&dev_number != -1)
      {
        priority_list[count] = dev_number;
        count++;
      }
    }
 
    tokens_ptr++;
  }
 
  g_strfreev(tokens);
 
  // terminate priority list with -1
  while(count < devs + 1) priority_list[count++] = -1;
 
  // opencl use can only be mandatory if at least one opencl device is given
  *mandatory = (priority_list[0] != -1) ? mnd : 0;
 
  dt_free(full);
}
 
// set device priorities according to config string
static void dt_opencl_update_priorities()
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited) return;
 
  // Priority parsing iterates over the list of available devices.
  // If !cl->inited, that means we have no available device, so empty list.
  // Exit early of face a segfault
  char *darkroom = dt_conf_get_string("opencl_devid_darkroom");
  char *preview = dt_conf_get_string("opencl_devid_preview");
  char *export = dt_conf_get_string("opencl_devid_export");
  char *thumbnail = dt_conf_get_string("opencl_devid_thumbnail");
 
  dt_opencl_priority_parse(cl, darkroom, cl->dev_priority_image, &cl->mandatory[0]);
  dt_opencl_priority_parse(cl, preview, cl->dev_priority_preview, &cl->mandatory[1]);
  dt_opencl_priority_parse(cl, export, cl->dev_priority_export, &cl->mandatory[2]);
  dt_opencl_priority_parse(cl, thumbnail, cl->dev_priority_thumbnail, &cl->mandatory[3]);
 
  dt_free(darkroom);
  dt_free(preview);
  dt_free(export);
  dt_free(thumbnail);
 
  dt_print_nts(DT_DEBUG_OPENCL, "[dt_opencl_update_priorities] these are your device priorities:\n");
  dt_print_nts(DT_DEBUG_OPENCL, "[dt_opencl_update_priorities] \tid |\t\tIMAGE\tPREVIEW\tEXPORT\tTHUMBS\n");
  for(int i = 0; i < cl->num_devs; i++)
    dt_print_nts(DT_DEBUG_OPENCL, "[dt_opencl_update_priorities]\t%i |\t\t%d\t%d\t%d\t%d\n",
                 i, cl->dev_priority_image[i],
                 cl->dev_priority_preview[i], cl->dev_priority_export[i], cl->dev_priority_thumbnail[i]);
  dt_print_nts(DT_DEBUG_OPENCL, "[dt_opencl_update_priorities] show if opencl use is mandatory for a given pixelpipe:\n");
  dt_print_nts(DT_DEBUG_OPENCL, "[dt_opencl_update_priorities] \t\tIMAGE\tPREVIEW\tEXPORT\tTHUMBS\n");
  dt_print_nts(DT_DEBUG_OPENCL, "[dt_opencl_update_priorities]\t\t%s\t%s\t%s\t%s\n", cl->mandatory[0] ? "yes" : "no",
             cl->mandatory[1] ? "yes" : "no", cl->mandatory[2] ? "yes" : "no", cl->mandatory[3] ? "yes" : "no");
}
 
int dt_opencl_lock_device(const int pipetype)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited) return -1;
 
 
  dt_pthread_mutex_lock(&cl->lock);
 
  size_t prio_size = sizeof(int) * (cl->num_devs + 1);
  int *priority = (int *)malloc(prio_size);
  int mandatory;
 
  switch(pipetype)
  {
    case DT_DEV_PIXELPIPE_FULL:
      memcpy(priority, cl->dev_priority_image, prio_size);
      mandatory = cl->mandatory[0];
      break;
    case DT_DEV_PIXELPIPE_PREVIEW:
      memcpy(priority, cl->dev_priority_preview, prio_size);
      mandatory = cl->mandatory[1];
      break;
    case DT_DEV_PIXELPIPE_EXPORT:
      memcpy(priority, cl->dev_priority_export, prio_size);
      mandatory = cl->mandatory[2];
      break;
    case DT_DEV_PIXELPIPE_THUMBNAIL:
      memcpy(priority, cl->dev_priority_thumbnail, prio_size);
      mandatory = cl->mandatory[3];
      break;
    default:
      dt_free(priority);
      mandatory = 0;
  }
 
  dt_pthread_mutex_unlock(&cl->lock);
 
  if(priority)
  {
    const int usec = 5000;
    const int nloop = MAX(0, dt_conf_get_int("opencl_mandatory_timeout"));
 
    // check for free opencl device repeatedly if mandatory is TRUE, else give up after first try
    for(int n = 0; n < nloop; n++)
    {
      const int *prio = priority;
 
      while(*prio != -1)
      {
        if(!dt_pthread_mutex_BAD_trylock(&cl->dev[*prio].lock))
        {
          int devid = *prio;
          dt_free(priority);
          return devid;
        }
        prio++;
      }
 
      if(!mandatory)
      {
        dt_free(priority);
        return -1;
      }
 
      dt_iop_nap(usec);
    }
    dt_print(DT_DEBUG_OPENCL, "[opencl_lock_device] reached opencl_mandatory_timeout trying to lock mandatory device, fallback to CPU\n");
  }
  else
  {
    // only a fallback if a new pipe type would be added and we forget to take care of it in opencl.c
    for(int try_dev = 0; try_dev < cl->num_devs; try_dev++)
    {
      // get first currently unused processor
      if(!dt_pthread_mutex_BAD_trylock(&cl->dev[try_dev].lock)) return try_dev;
    }
  }
 
  dt_free(priority);
 
  // no free GPU :(
  // use CPU processing, if no free device:
  return -1;
}
 
void dt_opencl_unlock_device(const int dev)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited) return;
  if(dev < 0 || dev >= cl->num_devs) return;
  dt_pthread_mutex_BAD_unlock(&cl->dev[dev].lock);
}
 
static FILE *fopen_stat(const char *filename, struct stat *st)
{
  FILE *f = g_fopen(filename, "rb");
  if(IS_NULL_PTR(f))
  {
    dt_print(DT_DEBUG_OPENCL, "[opencl_fopen_stat] could not open file `%s'!\n", filename);
    return NULL;
  }
  int fd = fileno(f);
  if(fstat(fd, st) < 0)
  {
    dt_print(DT_DEBUG_OPENCL, "[opencl_fopen_stat] could not stat file `%s'!\n", filename);
    return NULL;
  }
  return f;
}
 
 
void dt_opencl_md5sum(const char **files, char **md5sums)
{
  char kerneldir[PATH_MAX] = { 0 };
  char filename[PATH_MAX] = { 0 };
  dt_loc_get_kerneldir(kerneldir, sizeof(kerneldir));
 
  for(int n = 0; n < DT_OPENCL_MAX_INCLUDES; n++, files++, md5sums++)
  {
    if(!*files)
    {
      *md5sums = NULL;
      continue;
    }
 
    dt_concat_path_file(filename, kerneldir, *files);
 
    struct stat filestat;
    FILE *f = fopen_stat(filename, &filestat);
 
    if(IS_NULL_PTR(f))
    {
      dt_print(DT_DEBUG_OPENCL, "[opencl_md5sums] could not open file `%s'!\n", filename);
      *md5sums = NULL;
      continue;
    }
 
    size_t filesize = filestat.st_size;
    char *file = (char *)malloc(filesize);
 
    if(IS_NULL_PTR(file))
    {
      dt_print(DT_DEBUG_OPENCL, "[opencl_md5sums] could not allocate buffer for file `%s'!\n", filename);
      *md5sums = NULL;
      fclose(f);
      continue;
    }
 
    size_t rd = fread(file, sizeof(char), filesize, f);
    fclose(f);
 
    if(rd != filesize)
    {
      dt_free(file);
      dt_print(DT_DEBUG_OPENCL, "[opencl_md5sums] could not read all of file `%s'!\n", filename);
      *md5sums = NULL;
      continue;
    }
 
    *md5sums = g_compute_checksum_for_data(G_CHECKSUM_MD5, (guchar *)file, filesize);
 
    dt_free(file);
  }
}
 
int dt_opencl_load_program(const int dev, const int prog, const char *filename, const char *binname,
                           const char *cachedir, char *md5sum, char **includemd5, int *loaded_cached)
{
  cl_int err;
  dt_opencl_t *cl = darktable.opencl;
 
  struct stat filestat, cachedstat;
  *loaded_cached = 0;
 
  if(prog < 0 || prog >= DT_OPENCL_MAX_PROGRAMS)
  {
    dt_print(DT_DEBUG_OPENCL, "[opencl_load_source] invalid program number `%d' of file `%s'!\n", prog,
             filename);
    return 0;
  }
 
  if(cl->dev[dev].program_used[prog])
  {
    dt_print(DT_DEBUG_OPENCL,
             "[opencl_load_source] program number `%d' already in use when loading file `%s'!\n", prog,
             filename);
    return 0;
  }
 
  FILE *f = fopen_stat(filename, &filestat);
  if(IS_NULL_PTR(f)) return 0;
 
  size_t filesize = filestat.st_size;
  char *file = (char *)malloc(filesize + 2048);
  size_t rd = fread(file, sizeof(char), filesize, f);
  fclose(f);
  if(rd != filesize)
  {
    dt_free(file);
    dt_print(DT_DEBUG_OPENCL, "[opencl_load_source] could not read all of file `%s'!\n", filename);
    return 0;
  }
 
  char *start = file + filesize;
  char *end = start + 2048;
  size_t len;
 
  cl_device_id devid = cl->dev[dev].devid;
  (cl->dlocl->symbols->dt_clGetDeviceInfo)(devid, CL_DRIVER_VERSION, end - start, start, &len);
  start += len;
 
  cl_platform_id platform;
  (cl->dlocl->symbols->dt_clGetDeviceInfo)(devid, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &platform, NULL);
 
  (cl->dlocl->symbols->dt_clGetPlatformInfo)(platform, CL_PLATFORM_VERSION, end - start, start, &len);
  start += len;
 
  const char *options_md5 = cl->dev[dev].options_md5 ? cl->dev[dev].options_md5 : cl->dev[dev].options;
  len = g_strlcpy(start, options_md5, end - start);
  start += len;
 
  /* make sure that the md5sums of all the includes are applied as well */
  for(int n = 0; n < DT_OPENCL_MAX_INCLUDES; n++)
  {
    if(!includemd5[n]) continue;
    len = g_strlcpy(start, includemd5[n], end - start);
    start += len;
  }
 
  char *source_md5 = g_compute_checksum_for_data(G_CHECKSUM_MD5, (guchar *)file, start - file);
  g_strlcpy(md5sum, source_md5, 33);
  dt_free(source_md5);
 
  file[filesize] = '\0';
 
  char linkedfile[PATH_MAX] = { 0 };
  ssize_t linkedfile_len = 0;
 
#if defined(_WIN32)
  // No symlinks on Windows
  // Have to figure out the name using the filename + md5sum
  char dup[PATH_MAX] = { 0 };
  snprintf(dup, sizeof(dup), "%s.%s", binname, md5sum);
  FILE *cached = fopen_stat(dup, &cachedstat);
  g_strlcpy(linkedfile, md5sum, sizeof(linkedfile));
  linkedfile_len = strlen(md5sum);
#else
  FILE *cached = fopen_stat(binname, &cachedstat);
#endif
 
  if(cached)
  {
#if !defined(_WIN32)
    linkedfile_len = readlink(binname, linkedfile, sizeof(linkedfile) - 1);
#endif // !defined(_WIN32)
    if(linkedfile_len > 0)
    {
      linkedfile[linkedfile_len] = '\0';
 
      if(strncmp(linkedfile, md5sum, 33) == 0)
      {
        // md5sum matches, load cached binary
        size_t cached_filesize = cachedstat.st_size;
 
        unsigned char *cached_content = (unsigned char *)malloc(cached_filesize + 1);
        rd = fread(cached_content, sizeof(char), cached_filesize, cached);
        if(rd != cached_filesize)
        {
          dt_print(DT_DEBUG_OPENCL, "[opencl_load_program] could not read all of file '%s' MD5: %s!\n", binname, md5sum);
        }
        else
        {
          cl->dev[dev].program[prog] = (cl->dlocl->symbols->dt_clCreateProgramWithBinary)(
              cl->dev[dev].context, 1, &(cl->dev[dev].devid), &cached_filesize,
              (const unsigned char **)&cached_content, NULL, &err);
          if(err != CL_SUCCESS)
          {
            dt_print(DT_DEBUG_OPENCL,
                     "[opencl_load_program] could not load cached binary program from file '%s' MD5: '%s'! (%i)\n",
                     binname, md5sum, err);
          }
          else
          {
            cl->dev[dev].program_used[prog] = 1;
            *loaded_cached = 1;
          }
        }
        dt_free(cached_content);
      }
    }
    fclose(cached);
  }
 
 
  if(*loaded_cached == 0)
  {
    // if loading cached was unsuccessful for whatever reason,
    // try to remove cached binary & link
#if !defined(_WIN32)
    if(linkedfile_len > 0)
    {
      char link_dest[PATH_MAX] = { 0 };
      dt_concat_path_file(link_dest, cachedir, linkedfile);
      g_unlink(link_dest);
    }
    g_unlink(binname);
#else
    // delete the file which contains the MD5 name
    g_unlink(dup);
#endif 
 
    dt_print(DT_DEBUG_OPENCL,
             "[opencl_load_program] could not load cached binary program, trying to compile source\n");
 
    cl->dev[dev].program[prog] = (cl->dlocl->symbols->dt_clCreateProgramWithSource)(
        cl->dev[dev].context, 1, (const char **)&file, &filesize, &err);
    dt_free(file);
    if((err != CL_SUCCESS) || (cl->dev[dev].program[prog] == NULL))
    {
      dt_print(DT_DEBUG_OPENCL, "[opencl_load_source] could not create program from file `%s'! (%i)\n",
               filename, err);
      return 0;
    }
    else
    {
      cl->dev[dev].program_used[prog] = 1;
    }
  }
  else
  {
    dt_free(file);
    dt_vprint(DT_DEBUG_OPENCL, "[opencl_load_program] loaded cached binary program from file '%s' MD5: '%s' \n", binname, md5sum);
  }
 
  dt_vprint(DT_DEBUG_OPENCL, "[opencl_load_program] successfully loaded program from '%s' MD5: '%s'\n", filename, md5sum);
 
  return 1;
}
 
int dt_opencl_build_program(const int dev, const int prog, const char *binname, const char *cachedir,
                            char *md5sum, int loaded_cached)
{
  if(prog < 0 || prog >= DT_OPENCL_MAX_PROGRAMS) return -1;
  dt_opencl_t *cl = darktable.opencl;
  cl_program program = cl->dev[dev].program[prog];
  cl_int err = (cl->dlocl->symbols->dt_clBuildProgram)(program, 1, &(cl->dev[dev].devid), cl->dev[dev].options, 0, 0);
 
  if(err != CL_SUCCESS)
    dt_print(DT_DEBUG_OPENCL, "[opencl_build_program] could not build program: %i\n", err);
  else
    dt_vprint(DT_DEBUG_OPENCL, "[opencl_build_program] successfully built program\n");
 
  cl_build_status build_status;
  (cl->dlocl->symbols->dt_clGetProgramBuildInfo)(program, cl->dev[dev].devid, CL_PROGRAM_BUILD_STATUS,
                                                 sizeof(cl_build_status), &build_status, NULL);
  dt_vprint(DT_DEBUG_OPENCL, "[opencl_build_program] BUILD STATUS: %d\n", build_status);
 
  char *build_log;
  size_t ret_val_size;
  (cl->dlocl->symbols->dt_clGetProgramBuildInfo)(program, cl->dev[dev].devid, CL_PROGRAM_BUILD_LOG, 0, NULL,
                                                 &ret_val_size);
  if(ret_val_size != SIZE_MAX)
  {
    build_log = (char *)malloc(sizeof(char) * (ret_val_size + 1));
    if(build_log)
    {
      (cl->dlocl->symbols->dt_clGetProgramBuildInfo)(program, cl->dev[dev].devid, CL_PROGRAM_BUILD_LOG,
                                                     ret_val_size, build_log, NULL);
 
      build_log[ret_val_size] = '\0';
 
      dt_vprint(DT_DEBUG_OPENCL, "BUILD LOG:\n");
      dt_vprint(DT_DEBUG_OPENCL, "%s\n", build_log);
 
      dt_free(build_log);
    }
  }
 
  if(err != CL_SUCCESS)
    return err;
  else
  {
    if(!loaded_cached)
    {
      dt_vprint(DT_DEBUG_OPENCL, "[opencl_build_program] saving binary\n");
 
      cl_uint numdev = 0;
      err = (cl->dlocl->symbols->dt_clGetProgramInfo)(program, CL_PROGRAM_NUM_DEVICES, sizeof(cl_uint),
                                                      &numdev, NULL);
      if(err != CL_SUCCESS)
      {
        dt_print(DT_DEBUG_OPENCL, "[opencl_build_program] CL_PROGRAM_NUM_DEVICES failed: %i\n", err);
        return CL_SUCCESS;
      }
 
      cl_device_id *devices = malloc(sizeof(cl_device_id) * numdev);
      err = (cl->dlocl->symbols->dt_clGetProgramInfo)(program, CL_PROGRAM_DEVICES,
                                                      sizeof(cl_device_id) * numdev, devices, NULL);
      if(err != CL_SUCCESS)
      {
        dt_print(DT_DEBUG_OPENCL, "[opencl_build_program] CL_PROGRAM_DEVICES failed: %i\n", err);
        dt_free(devices);
        return CL_SUCCESS;
      }
 
      size_t *binary_sizes = malloc(sizeof(size_t) * numdev);
      err = (cl->dlocl->symbols->dt_clGetProgramInfo)(program, CL_PROGRAM_BINARY_SIZES,
                                                      sizeof(size_t) * numdev, binary_sizes, NULL);
      if(err != CL_SUCCESS)
      {
        dt_print(DT_DEBUG_OPENCL, "[opencl_build_program] CL_PROGRAM_BINARY_SIZES failed: %i\n", err);
        dt_free(binary_sizes);
        dt_free(devices);
        return CL_SUCCESS;
      }
 
      unsigned char **binaries = malloc(sizeof(unsigned char *) * numdev);
      for(int i = 0; i < numdev; i++) binaries[i] = (unsigned char *)malloc(binary_sizes[i]);
      err = (cl->dlocl->symbols->dt_clGetProgramInfo)(program, CL_PROGRAM_BINARIES,
                                                      sizeof(unsigned char *) * numdev, binaries, NULL);
      if(err != CL_SUCCESS)
      {
        dt_print(DT_DEBUG_OPENCL, "[opencl_build_program] CL_PROGRAM_BINARIES failed: %i\n", err);
        goto ret;
      }
 
      for(int i = 0; i < numdev; i++)
        if(cl->dev[dev].devid == devices[i])
        {
          // save opencl compiled binary as md5sum-named file
          char link_dest[PATH_MAX] = { 0 };
          snprintf(link_dest, sizeof(link_dest), "%s" G_DIR_SEPARATOR_S "%s", cachedir, md5sum);
          FILE *f = g_fopen(link_dest, "wb");
          if(IS_NULL_PTR(f)) goto ret;
          size_t bytes_written = fwrite(binaries[i], sizeof(char), binary_sizes[i], f);
          if(bytes_written != binary_sizes[i]) goto ret;
          fclose(f);
 
          // create link (e.g. basic.cl.bin -> f1430102c53867c162bb60af6c163328)
          char cwd[PATH_MAX] = { 0 };
          if(!getcwd(cwd, sizeof(cwd))) goto ret;
          if(chdir(cachedir) != 0) goto ret;
          char dup[PATH_MAX] = { 0 };
          g_strlcpy(dup, binname, sizeof(dup));
          char *bname = basename(dup);
#if defined(_WIN32)
          //CreateSymbolicLink in Windows requires admin privileges, which we don't want/need
          //store has using a simple filerename
          char finalfilename[PATH_MAX] = { 0 };
          snprintf(finalfilename, sizeof(finalfilename), "%s" G_DIR_SEPARATOR_S "%s.%s", cachedir, bname, md5sum);
          rename(link_dest, finalfilename);
#else
          if(symlink(md5sum, bname) != 0) goto ret;
#endif 
          if(chdir(cwd) != 0) goto ret;
        }
 
    ret:
      for(int i = 0; i < numdev; i++) dt_free(binaries[i]);
      dt_free(binaries);
      dt_free(binary_sizes);
      dt_free(devices);
    }
    return CL_SUCCESS;
  }
}
 
int dt_opencl_create_kernel(const int prog, const char *name)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited) return -1;
  if(prog < 0 || prog >= DT_OPENCL_MAX_PROGRAMS) return -1;
  dt_pthread_mutex_lock(&cl->lock);
  int k = 0;
  for(int dev = 0; dev < cl->num_devs; dev++)
  {
    cl_int err;
    for(; k < DT_OPENCL_MAX_KERNELS; k++)
      if(!cl->dev[dev].kernel_used[k])
      {
        cl->dev[dev].kernel_used[k] = 1;
        cl->dev[dev].kernel[k]
            = (cl->dlocl->symbols->dt_clCreateKernel)(cl->dev[dev].program[prog], name, &err);
        if(err != CL_SUCCESS)
        {
          dt_print(DT_DEBUG_OPENCL, "[opencl_create_kernel] could not create kernel `%s'! (%i)\n", name, err);
          cl->dev[dev].kernel_used[k] = 0;
          goto error;
        }
        else
          break;
      }
    if(k < DT_OPENCL_MAX_KERNELS)
    {
      dt_vprint(DT_DEBUG_OPENCL, "[opencl_create_kernel] successfully loaded kernel `%s' (%d) for device %d\n",
               name, k, dev);
    }
    else
    {
      dt_print(DT_DEBUG_OPENCL, "[opencl_create_kernel] too many kernels! can't create kernel `%s'\n", name);
      goto error;
    }
  }
  dt_pthread_mutex_unlock(&cl->lock);
  return k;
error:
  dt_pthread_mutex_unlock(&cl->lock);
  return -1;
}
 
void dt_opencl_free_kernel(const int kernel)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited) return;
  if(kernel < 0 || kernel >= DT_OPENCL_MAX_KERNELS) return;
  dt_pthread_mutex_lock(&cl->lock);
  for(int dev = 0; dev < cl->num_devs; dev++)
  {
    cl->dev[dev].kernel_used[kernel] = 0;
    (cl->dlocl->symbols->dt_clReleaseKernel)(cl->dev[dev].kernel[kernel]);
  }
  dt_pthread_mutex_unlock(&cl->lock);
}
 
int dt_opencl_get_max_work_item_sizes(const int dev, size_t *sizes)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited || dev < 0) return -1;
  return (cl->dlocl->symbols->dt_clGetDeviceInfo)(cl->dev[dev].devid, CL_DEVICE_MAX_WORK_ITEM_SIZES,
                                                  sizeof(size_t) * 3, sizes, NULL);
}
 
int dt_opencl_get_work_group_limits(const int dev, size_t *sizes, size_t *workgroupsize,
                                    unsigned long *localmemsize)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited || dev < 0) return -1;
  cl_ulong lmemsize;
  cl_int err = (cl->dlocl->symbols->dt_clGetDeviceInfo)(cl->dev[dev].devid, CL_DEVICE_LOCAL_MEM_SIZE,
                                                 sizeof(cl_ulong), &lmemsize, NULL);
  if(err != CL_SUCCESS) return err;
 
  *localmemsize = lmemsize;
 
  err = (cl->dlocl->symbols->dt_clGetDeviceInfo)(cl->dev[dev].devid, CL_DEVICE_MAX_WORK_GROUP_SIZE,
                                                 sizeof(size_t), workgroupsize, NULL);
  if(err != CL_SUCCESS) return err;
 
  return dt_opencl_get_max_work_item_sizes(dev, sizes);
}
 
 
int dt_opencl_get_kernel_work_group_size(const int dev, const int kernel, size_t *kernelworkgroupsize)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited || dev < 0) return -1;
  if(kernel < 0 || kernel >= DT_OPENCL_MAX_KERNELS) return -1;
 
  return (cl->dlocl->symbols->dt_clGetKernelWorkGroupInfo)(cl->dev[dev].kernel[kernel], cl->dev[dev].devid,
                                                           CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t),
                                                           kernelworkgroupsize, NULL);
}
 
 
int dt_opencl_set_kernel_arg(const int dev, const int kernel, const int num, const size_t size,
                             const void *arg)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited || dev < 0) return -1;
  if(kernel < 0 || kernel >= DT_OPENCL_MAX_KERNELS) return -1;
  return (cl->dlocl->symbols->dt_clSetKernelArg)(cl->dev[dev].kernel[kernel], num, size, arg);
}
 
int dt_opencl_enqueue_kernel_2d(const int dev, const int kernel, const size_t *sizes)
{
  return dt_opencl_enqueue_kernel_2d_with_local(dev, kernel, sizes, NULL);
}
 
 
int dt_opencl_enqueue_kernel_2d_with_local(const int dev, const int kernel, const size_t *sizes,
                                           const size_t *local)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited || dev < 0) return -1;
  if(kernel < 0 || kernel >= DT_OPENCL_MAX_KERNELS) return -1;
 
  char buf[256];
  buf[0] = '\0';
  if(darktable.unmuted & DT_DEBUG_OPENCL)
    (cl->dlocl->symbols->dt_clGetKernelInfo)(cl->dev[dev].kernel[kernel], CL_KERNEL_FUNCTION_NAME, 256, buf, NULL);
  cl_event *eventp = dt_opencl_events_get_slot(dev, buf);
  cl_int err = (cl->dlocl->symbols->dt_clEnqueueNDRangeKernel)(cl->dev[dev].cmd_queue, cl->dev[dev].kernel[kernel],
                                                        2, NULL, sizes, local, 0, NULL, eventp);
 
  if(err != CL_SUCCESS)
    dt_print(DT_DEBUG_OPENCL, "[dt_opencl_enqueue_kernel_2d_with_local] kernel %i (%s) on device %d: %i\n", kernel, buf, dev, err);
 
  return err;
}
 
int dt_opencl_copy_device_to_host(const int devid, void *host, void *device, const int width,
                                  const int height, const int bpp)
{
  return dt_opencl_read_host_from_device(devid, host, device, width, height, bpp);
}
 
int dt_opencl_read_host_from_device(const int devid, void *host, void *device, const int width,
                                    const int height, const int bpp)
{
  return dt_opencl_read_host_from_device_rowpitch(devid, host, device, width, height, bpp * width);
}
 
int dt_opencl_read_host_from_device_rowpitch(const int devid, void *host, void *device, const int width,
                                             const int height, const int rowpitch)
{
  if(!darktable.opencl->inited || devid < 0) return -1;
  const size_t origin[] = { 0, 0, 0 };
  const size_t region[] = { width, height, 1 };
  // blocking.
  return dt_opencl_read_host_from_device_raw(devid, host, device, origin, region, rowpitch, CL_TRUE);
}
 
int dt_opencl_read_host_from_device_non_blocking(const int devid, void *host, void *device, const int width,
                                                 const int height, const int bpp)
{
  return dt_opencl_read_host_from_device_rowpitch_non_blocking(devid, host, device, width, height,
                                                               bpp * width);
}
 
int dt_opencl_read_host_from_device_rowpitch_non_blocking(const int devid, void *host, void *device,
                                                          const int width, const int height,
                                                          const int rowpitch)
{
  if(!darktable.opencl->inited || devid < 0) return -1;
  const size_t origin[] = { 0, 0, 0 };
  const size_t region[] = { width, height, 1 };
  // non-blocking.
  return dt_opencl_read_host_from_device_raw(devid, host, device, origin, region, rowpitch, CL_FALSE);
}
 
 
int dt_opencl_read_host_from_device_raw(const int devid, void *host, void *device, const size_t *origin,
                                        const size_t *region, const int rowpitch, const int blocking)
{
  if(!darktable.opencl->inited) return -1;
 
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Read Image (from device to host)]");
 
  return (darktable.opencl->dlocl->symbols->dt_clEnqueueReadImage)(darktable.opencl->dev[devid].cmd_queue,
                                                                   device, blocking ? CL_TRUE : CL_FALSE, origin, region, rowpitch,
                                                                   0, host, 0, NULL, eventp);
}
 
int dt_opencl_write_host_to_device(const int devid, void *host, void *device, const int width,
                                   const int height, const int bpp)
{
  return dt_opencl_write_host_to_device_rowpitch(devid, host, device, width, height, width * bpp);
}
 
int dt_opencl_write_host_to_device_rowpitch(const int devid, void *host, void *device, const int width,
                                            const int height, const int rowpitch)
{
  if(!darktable.opencl->inited || devid < 0) return -1;
  const size_t origin[] = { 0, 0, 0 };
  const size_t region[] = { width, height, 1 };
  // blocking.
  return dt_opencl_write_host_to_device_raw(devid, host, device, origin, region, rowpitch, CL_TRUE);
}
 
int dt_opencl_write_host_to_device_non_blocking(const int devid, void *host, void *device, const int width,
                                                const int height, const int bpp)
{
  return dt_opencl_write_host_to_device_rowpitch_non_blocking(devid, host, device, width, height, width * bpp);
}
 
int dt_opencl_write_host_to_device_rowpitch_non_blocking(const int devid, void *host, void *device,
                                                         const int width, const int height,
                                                         const int rowpitch)
{
  if(!darktable.opencl->inited || devid < 0) return -1;
  const size_t origin[] = { 0, 0, 0 };
  const size_t region[] = { width, height, 1 };
  // non-blocking.
  return dt_opencl_write_host_to_device_raw(devid, host, device, origin, region, rowpitch, CL_FALSE);
}
 
int dt_opencl_write_host_to_device_raw(const int devid, const void *host, void *device, const size_t *origin,
                                       const size_t *region, const int rowpitch, const int blocking)
{
  if(!darktable.opencl->inited) return -1;
 
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Write Image (from host to device)]");
 
  return (darktable.opencl->dlocl->symbols->dt_clEnqueueWriteImage)(darktable.opencl->dev[devid].cmd_queue,
                                                                    device, blocking ? CL_TRUE : CL_FALSE, origin, region,
                                                                    rowpitch, 0, host, 0, NULL, eventp);
}
 
int dt_opencl_enqueue_copy_image(const int devid, cl_mem src, cl_mem dst, size_t *orig_src, size_t *orig_dst,
                                 size_t *region)
{
  if(!darktable.opencl->inited || devid < 0) return -1;
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Copy Image (on device)]");
  cl_int err = (darktable.opencl->dlocl->symbols->dt_clEnqueueCopyImage)(
      darktable.opencl->dev[devid].cmd_queue, src, dst, orig_src, orig_dst, region, 0, NULL, eventp);
  if(err != CL_SUCCESS) dt_print(DT_DEBUG_OPENCL, "[opencl copy_image] could not copy image on device %d: %i\n", devid, err);
  return err;
}
 
int dt_opencl_enqueue_copy_image_to_buffer(const int devid, cl_mem src_image, cl_mem dst_buffer,
                                           size_t *origin, size_t *region, size_t offset)
{
  if(!darktable.opencl->inited) return -1;
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Copy Image to Buffer (on device)]");
  cl_int err = (darktable.opencl->dlocl->symbols->dt_clEnqueueCopyImageToBuffer)(
      darktable.opencl->dev[devid].cmd_queue, src_image, dst_buffer, origin, region, offset, 0, NULL, eventp);
  if(err != CL_SUCCESS)
    dt_print(DT_DEBUG_OPENCL, "[opencl copy_image_to_buffer] could not copy image on device %d: %i\n", devid, err);
  return err;
}
 
int dt_opencl_enqueue_copy_buffer_to_image(const int devid, cl_mem src_buffer, cl_mem dst_image,
                                           size_t offset, size_t *origin, size_t *region)
{
  if(!darktable.opencl->inited) return -1;
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Copy Buffer to Image (on device)]");
  cl_int err = (darktable.opencl->dlocl->symbols->dt_clEnqueueCopyBufferToImage)(
      darktable.opencl->dev[devid].cmd_queue, src_buffer, dst_image, offset, origin, region, 0, NULL, eventp);
  if(err != CL_SUCCESS)
    dt_print(DT_DEBUG_OPENCL, "[opencl copy_buffer_to_image] could not copy buffer on device %d: %i\n", devid, err);
  return err;
}
 
int dt_opencl_enqueue_copy_buffer_to_buffer(const int devid, cl_mem src_buffer, cl_mem dst_buffer,
                                            size_t srcoffset, size_t dstoffset, size_t size)
{
  if(!darktable.opencl->inited) return -1;
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Copy Buffer to Buffer (on device)]");
  cl_int err = (darktable.opencl->dlocl->symbols->dt_clEnqueueCopyBuffer)(darktable.opencl->dev[devid].cmd_queue,
                                                                   src_buffer, dst_buffer, srcoffset,
                                                                   dstoffset, size, 0, NULL, eventp);
  if(err != CL_SUCCESS)
    dt_print(DT_DEBUG_OPENCL, "[opencl copy_buffer_to_buffer] could not copy buffer on device %d: %i\n", devid, err);
  return err;
}
 
int dt_opencl_read_buffer_from_device(const int devid, void *host, void *device, const size_t offset,
                                      const size_t size, const int blocking)
{
  if(!darktable.opencl->inited) return -1;
 
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Read Buffer (from device to host)]");
 
  return (darktable.opencl->dlocl->symbols->dt_clEnqueueReadBuffer)(
      darktable.opencl->dev[devid].cmd_queue, device, blocking ? CL_TRUE : CL_FALSE, offset, size, host, 0, NULL, eventp);
}
 
int dt_opencl_write_buffer_to_device(const int devid, void *host, void *device, const size_t offset,
                                     const size_t size, const int blocking)
{
  if(!darktable.opencl->inited) return -1;
 
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Write Buffer (from host to device)]");
 
  return (darktable.opencl->dlocl->symbols->dt_clEnqueueWriteBuffer)(
      darktable.opencl->dev[devid].cmd_queue, device, blocking ? CL_TRUE : CL_FALSE, offset, size, host, 0, NULL, eventp);
}
 
 
void *dt_opencl_copy_host_to_device_constant(const int devid, const size_t size, void *host)
{
  if(!darktable.opencl->inited || devid < 0) return NULL;
  cl_int err;
  cl_mem dev = (darktable.opencl->dlocl->symbols->dt_clCreateBuffer)(
      darktable.opencl->dev[devid].context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, size, host, &err);
  if(err != CL_SUCCESS)
    dt_print(DT_DEBUG_OPENCL,
             "[opencl copy_host_to_device_constant] could not alloc buffer on device %d: %i\n", devid, err);
 
  if(err == CL_SUCCESS) dt_opencl_memory_statistics(devid, dev, size, OPENCL_MEMORY_ADD);
 
  return dev;
}
 
void *dt_opencl_copy_host_to_device(const int devid, void *host, const int width, const int height,
                                    const int bpp)
{
  return dt_opencl_copy_host_to_device_rowpitch(devid, host, width, height, bpp, 0);
}
 
void *dt_opencl_copy_host_to_device_rowpitch(const int devid, void *host, const int width, const int height,
                                             const int bpp, const int rowpitch)
{
  if(!darktable.opencl->inited || devid < 0) return NULL;
  cl_int err;
  cl_image_format fmt;
  // guess pixel format from bytes per pixel
  if(bpp == 4 * sizeof(float))
    fmt = (cl_image_format){ CL_RGBA, CL_FLOAT };
  else if(bpp == sizeof(float))
    fmt = (cl_image_format){ CL_R, CL_FLOAT };
  else if(bpp == sizeof(uint16_t))
    fmt = (cl_image_format){ CL_R, CL_UNSIGNED_INT16 };
  else
    return NULL;
 
  // TODO: if fmt = uint16_t, blow up to 4xuint16_t and copy manually!
  cl_mem dev = (darktable.opencl->dlocl->symbols->dt_clCreateImage2D)(
      darktable.opencl->dev[devid].context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, &fmt, width, height,
      rowpitch, host, &err);
  if(err != CL_SUCCESS)
    dt_print(DT_DEBUG_OPENCL,
             "[opencl copy_host_to_device] could not alloc/copy img buffer on device %d: %i\n", devid, err);
 
  if(err == CL_SUCCESS)
  {
    const size_t bytes = (size_t)(rowpitch ? rowpitch : width * bpp) * height;
    dt_opencl_memory_statistics(devid, dev, bytes, OPENCL_MEMORY_ADD);
  }
 
  return dev;
}
 
 
void dt_opencl_release_mem_object(cl_mem mem)
{
  if(!darktable.opencl->inited) return;
 
  // the OpenCL specs are not absolutely clear if clReleaseMemObject(NULL) is a no-op. we take care of the
  // case in a centralized way at this place
  if(IS_NULL_PTR(mem)) return;
 
  dt_opencl_memory_statistics(-1, mem, 0, OPENCL_MEMORY_SUB);
 
  (darktable.opencl->dlocl->symbols->dt_clReleaseMemObject)(mem);
}
 
void *dt_opencl_map_buffer(const int devid, cl_mem buffer, const int blocking, const int flags, size_t offset,
                           size_t size)
{
  if(!darktable.opencl->inited) return NULL;
  cl_int err;
  void *ptr;
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Map Buffer]");
  ptr = (darktable.opencl->dlocl->symbols->dt_clEnqueueMapBuffer)(
      darktable.opencl->dev[devid].cmd_queue, buffer, blocking ? CL_TRUE : CL_FALSE, flags, offset, size, 0, NULL, eventp, &err);
  if(err != CL_SUCCESS) dt_print(DT_DEBUG_OPENCL, "[opencl map buffer] could not map buffer on device %d: %i\n", devid, err);
  return ptr;
}
 
 
void *dt_opencl_map_image(const int devid, cl_mem buffer, const int blocking, const int flags, size_t width, size_t height, int bpp)
{
  if(!darktable.opencl->inited) return NULL;
  cl_int err;
  void *ptr;
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Map Image 2D]");
  size_t origin[3] = {0, 0, 0};
  size_t region[3] = {width, height, 1};
  size_t mapped_row_pitch;
 
  ptr = (darktable.opencl->dlocl->symbols->dt_clEnqueueMapImage)(
      darktable.opencl->dev[devid].cmd_queue, buffer, blocking ? CL_TRUE : CL_FALSE, flags, origin, region,
      &mapped_row_pitch, NULL, 0, NULL, eventp, &err);
 
  if(err != CL_SUCCESS)
    dt_print(DT_DEBUG_OPENCL, "[opencl map buffer] could not map image on device %d: %i\n", devid, err);
  return ptr;
}
 
 
int dt_opencl_unmap_mem_object(const int devid, cl_mem mem_object, void *mapped_ptr)
{
  if(!darktable.opencl->inited) return -1;
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Unmap Mem Object]");
  cl_int err = (darktable.opencl->dlocl->symbols->dt_clEnqueueUnmapMemObject)(
      darktable.opencl->dev[devid].cmd_queue, mem_object, mapped_ptr, 0, NULL, eventp);
  if(err != CL_SUCCESS)
    dt_print(DT_DEBUG_OPENCL, "[opencl unmap mem object] could not unmap mem object on device %d: %i\n", devid, err);
  return err;
}
 
static inline void *_dt_opencl_alloc_image2d(const int devid, const int width, const int height,
                                             const size_t bytes, const cl_mem_flags flags,
                                             const cl_image_format fmt, void *host,
                                             const char *const context)
{
  if(!darktable.opencl->inited || devid < 0) return NULL;
  cl_int err;
  cl_mem dev = NULL;
  for(int attempt = 0; attempt < 2; attempt++)
  {
    dev = (darktable.opencl->dlocl->symbols->dt_clCreateImage2D)(darktable.opencl->dev[devid].context, flags,
                                                                  &fmt, width, height, 0, host, &err);
    if(err == CL_SUCCESS) break;
    if(attempt == 0 && (err == CL_MEM_OBJECT_ALLOCATION_FAILURE || err == CL_OUT_OF_RESOURCES))
    {
      dt_print(DT_DEBUG_OPENCL,
               "[opencl %s] out of memory on device %d, flushing cached pinned buffers and retrying\n",
               context, devid);
      dt_dev_pixelpipe_cache_flush_clmem(darktable.pixelpipe_cache, devid);
      continue;
    }
    break;
  }
 
  if(err != CL_SUCCESS)
    dt_print(DT_DEBUG_OPENCL, "[opencl %s] could not alloc img buffer on device %d: %i\n", context, devid, err);
 
  if(err == CL_SUCCESS) dt_opencl_memory_statistics(devid, dev, bytes, OPENCL_MEMORY_ADD);
  return dev;
}
 
void *dt_opencl_alloc_device(const int devid, const int width, const int height, const int bpp)
{
  const int effective_bpp = DT_OPENCL_BPP_DECODE(bpp);
  const gboolean rgba8 = DT_OPENCL_BPP_IS_RGBA8(bpp);
  cl_image_format fmt;
  // guess pixel format from bytes per pixel (+ optional format tag for ambiguous 4-byte formats)
  if(rgba8 && effective_bpp == 4 * sizeof(uint8_t))
    fmt = (cl_image_format){ CL_RGBA, CL_UNSIGNED_INT8 };
  else if(effective_bpp == 4 * sizeof(float))
    fmt = (cl_image_format){ CL_RGBA, CL_FLOAT };
  else if(effective_bpp == sizeof(float))
    fmt = (cl_image_format){ CL_R, CL_FLOAT };
  else if(effective_bpp == sizeof(uint16_t))
    fmt = (cl_image_format){ CL_R, CL_UNSIGNED_INT16 };
  else if(effective_bpp == sizeof(uint8_t))
    fmt = (cl_image_format){ CL_R, CL_UNSIGNED_INT8 };
  else
    return NULL;
 
  const size_t bytes = (size_t)width * height * effective_bpp;
  return _dt_opencl_alloc_image2d(devid, width, height, bytes, CL_MEM_READ_WRITE, fmt, NULL, "alloc_device");
}
 
void *dt_opencl_alloc_device_use_host_pointer(const int devid, const int width, const int height,
                                              const int bpp, void *host, const int flags)
{
  const int effective_bpp = DT_OPENCL_BPP_DECODE(bpp);
  const gboolean rgba8 = DT_OPENCL_BPP_IS_RGBA8(bpp);
  cl_image_format fmt;
  // guess pixel format from bytes per pixel (+ optional format tag for ambiguous 4-byte formats)
  if(rgba8 && effective_bpp == 4 * sizeof(uint8_t))
    fmt = (cl_image_format){ CL_RGBA, CL_UNSIGNED_INT8 };
  else if(effective_bpp == 4 * sizeof(float))
    fmt = (cl_image_format){ CL_RGBA, CL_FLOAT };
  else if(effective_bpp == sizeof(float))
    fmt = (cl_image_format){ CL_R, CL_FLOAT };
  else if(effective_bpp == sizeof(uint16_t))
    fmt = (cl_image_format){ CL_R, CL_UNSIGNED_INT16 };
  else
    return NULL;
 
  const size_t bytes = (size_t)width * height * effective_bpp;
  return _dt_opencl_alloc_image2d(devid, width, height, bytes, flags, fmt, host,
                                  "alloc_device_use_host_pointer");
}
 
void *dt_opencl_alloc_device_buffer_with_flags(const int devid, const size_t size, const int flags, void *host_ptr)
{
  if(!darktable.opencl->inited) return NULL;
  cl_int err;
  cl_mem buf = NULL;
  for(int attempt = 0; attempt < 2; attempt++)
  {
    buf = (darktable.opencl->dlocl->symbols->dt_clCreateBuffer)(darktable.opencl->dev[devid].context,
                                                               flags, size, host_ptr, &err);
    if(err == CL_SUCCESS) break;
    if(attempt == 0 && (err == CL_MEM_OBJECT_ALLOCATION_FAILURE || err == CL_OUT_OF_RESOURCES))
    {
      dt_print(DT_DEBUG_OPENCL,
               "[opencl alloc_device_buffer] out of memory on device %d, flushing cached pinned buffers and retrying\n",
               devid);
      dt_dev_pixelpipe_cache_flush_clmem(darktable.pixelpipe_cache, devid);
      continue;
    }
    break;
  }
  if(err != CL_SUCCESS)
    dt_print(DT_DEBUG_OPENCL, "[opencl alloc_device_buffer] could not alloc buffer on device %d: %d\n", devid,
             err);
 
  if(err == CL_SUCCESS) dt_opencl_memory_statistics(devid, buf, size, OPENCL_MEMORY_ADD);
 
  return buf;
}
 
 
void *dt_opencl_alloc_device_buffer(const int devid, const size_t size)
{
  return dt_opencl_alloc_device_buffer_with_flags(devid, size, CL_MEM_READ_WRITE, NULL);
}
 
 
size_t dt_opencl_get_mem_object_size(cl_mem mem)
{
  size_t size;
  if(IS_NULL_PTR(mem)) return 0;
 
  cl_int err = (darktable.opencl->dlocl->symbols->dt_clGetMemObjectInfo)(mem, CL_MEM_SIZE, sizeof(size), &size, NULL);
 
  return (err == CL_SUCCESS) ? size : 0;
}
 
int dt_opencl_get_mem_context_id(cl_mem mem)
{
  cl_context context;
  if(IS_NULL_PTR(mem)) return -1;
 
  cl_int err = (darktable.opencl->dlocl->symbols->dt_clGetMemObjectInfo)(mem, CL_MEM_CONTEXT, sizeof(context), &context, NULL);
  if(err != CL_SUCCESS)
    return -1;
 
  for(int devid = 0; devid < darktable.opencl->num_devs; devid++)
  {
    if(darktable.opencl->dev[devid].context == context)
      return devid;
  }
 
  return -1;
}
 
cl_mem_flags dt_opencl_get_mem_flags(cl_mem mem)
{
  if(!darktable.opencl->inited || IS_NULL_PTR(mem)) return 0;
  cl_mem_flags flags = 0;
  cl_int err = (darktable.opencl->dlocl->symbols->dt_clGetMemObjectInfo)(mem, CL_MEM_FLAGS, sizeof(flags), &flags, NULL);
  if(err != CL_SUCCESS) return 0;
  return flags;
}
 
int dt_opencl_get_image_width(cl_mem mem)
{
  size_t size;
  if(IS_NULL_PTR(mem)) return 0;
 
  cl_int err = (darktable.opencl->dlocl->symbols->dt_clGetImageInfo)(mem, CL_IMAGE_WIDTH, sizeof(size), &size, NULL);
  if(size > INT_MAX) size = 0;
 
  return (err == CL_SUCCESS) ? (int)size : 0;
}
 
int dt_opencl_get_image_height(cl_mem mem)
{
  size_t size;
  if(IS_NULL_PTR(mem)) return 0;
 
  cl_int err = (darktable.opencl->dlocl->symbols->dt_clGetImageInfo)(mem, CL_IMAGE_HEIGHT, sizeof(size), &size, NULL);
  if(size > INT_MAX) size = 0;
 
  return (err == CL_SUCCESS) ? (int)size : 0;
}
 
int dt_opencl_get_image_element_size(cl_mem mem)
{
  size_t size;
  if(IS_NULL_PTR(mem)) return 0;
 
  cl_int err = (darktable.opencl->dlocl->symbols->dt_clGetImageInfo)(mem, CL_IMAGE_ELEMENT_SIZE, sizeof(size), &size,
                                                              NULL);
  if(size > INT_MAX) size = 0;
 
  return (err == CL_SUCCESS) ? (int)size : 0;
}
 
typedef struct dt_opencl_mem_record_t
{
  int devid;
  size_t size;
} dt_opencl_mem_record_t;
 
void dt_opencl_memory_statistics(int devid, cl_mem mem, size_t size, dt_opencl_memory_t action)
{
  if(IS_NULL_PTR(mem)) return;
 
  if(action == OPENCL_MEMORY_ADD)
  {
    // devid and size are known at allocation time -- record them so the release
    // path can undo the exact same amount without querying the driver.
    if(devid < 0) return;
    dt_opencl_mem_record_t *rec = (dt_opencl_mem_record_t *)malloc(sizeof(dt_opencl_mem_record_t));
    rec->devid = devid;
    rec->size = size;
    dt_pthread_mutex_lock(&darktable.opencl->mem_sizes_lock);
    // g_hash_table_insert frees the previous value (g_free) if the key already
    // exists, so re-inserting the same cl_mem pointer never leaks.
    g_hash_table_insert(darktable.opencl->mem_sizes, mem, rec);
    dt_pthread_mutex_unlock(&darktable.opencl->mem_sizes_lock);
  }
  else
  {
    // Look up what we recorded on ADD. Never ask the driver about the object here
    // either: on some Windows drivers clGetMemObjectInfo faults under vRAM
    // pressure (issues #130221119 / #130557353), aborting the process.
    dt_pthread_mutex_lock(&darktable.opencl->mem_sizes_lock);
    dt_opencl_mem_record_t *rec = (dt_opencl_mem_record_t *)g_hash_table_lookup(darktable.opencl->mem_sizes, mem);
    if(rec)
    {
      devid = rec->devid;
      size = rec->size;
      g_hash_table_remove(darktable.opencl->mem_sizes, mem);
    }
    else
    {
      // Untracked object (allocated before this bookkeeping, or via a path that
      // did not register). Nothing reliable to subtract; leave the counter be.
      devid = -1;
    }
    dt_pthread_mutex_unlock(&darktable.opencl->mem_sizes_lock);
  }
 
  if(devid < 0)
    return;
 
  if(action == OPENCL_MEMORY_ADD)
    darktable.opencl->dev[devid].memory_in_use += size;
  else
    darktable.opencl->dev[devid].memory_in_use =
      (darktable.opencl->dev[devid].memory_in_use > size)
        ? (darktable.opencl->dev[devid].memory_in_use - size)
        : 0;
 
  darktable.opencl->dev[devid].peak_memory = MAX(darktable.opencl->dev[devid].peak_memory,
                                                 darktable.opencl->dev[devid].memory_in_use);
 
  if((darktable.unmuted & DT_DEBUG_MEMORY) && (darktable.unmuted & DT_DEBUG_OPENCL))
    dt_print(DT_DEBUG_OPENCL,
              "[opencl memory] device %d: %" G_GSIZE_FORMAT " bytes (%.1f MB) in use\n", devid, darktable.opencl->dev[devid].memory_in_use,
                                      (float)darktable.opencl->dev[devid].memory_in_use/(1024*1024));
}
 
void dt_opencl_check_tuning(const int devid)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited || devid < 0) return;
 
  // Apply the headroom read from this device configuration. Older configs without
  // a per-device key are migrated from the global default during device init.
  size_t headroom = cl->dev[devid].forced_headroom;
 
  cl->dev[devid].used_available = MAX(0ul, cl->dev[devid].max_global_mem - headroom * 1024 * 1024);
 
  dt_print(DT_DEBUG_OPENCL | DT_DEBUG_MEMORY,
      "[dt_opencl_check_tuning] use %" G_GSIZE_FORMAT " MiB on device `%s' id=%i\n",
      cl->dev[devid].used_available / (1024 * 1024),
      cl->dev[devid].name, devid);
}
 
cl_ulong dt_opencl_get_device_available(const int devid)
{
  if(!darktable.opencl->inited || devid < 0) return 0;
  const cl_ulong limit = darktable.opencl->dev[devid].used_available;
  const size_t in_use = darktable.opencl->dev[devid].memory_in_use;
  return (limit > in_use) ? (limit - in_use) : 0;
}
 
static cl_ulong _opencl_get_device_memalloc(const int devid)
{
  return darktable.opencl->dev[devid].max_mem_alloc;
}
 
cl_ulong dt_opencl_get_device_memalloc(const int devid)
{
  if(!darktable.opencl->inited || devid < 0) return 0;
  return _opencl_get_device_memalloc(devid);
}
 
dt_opencl_fit_reason_t dt_opencl_image_fits_device_reason(const int devid, const size_t width,
                                const size_t height, const unsigned bpp, const float factor,
                                const size_t overhead, size_t *needed, size_t *limit)
{
  size_t n = 0, l = 0;
  dt_opencl_fit_reason_t reason = DT_OPENCL_FIT_OK;
 
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited || devid < 0)
  {
    reason = DT_OPENCL_FIT_UNINITED;
  }
  else
  {
    const size_t required = width * height * bpp;
    const size_t total = (size_t)ceilf((float)required * factor) + overhead;
 
    if(cl->dev[devid].max_image_width < width || cl->dev[devid].max_image_height < height)
    {
      // dimension limit: compared quantities are pixel counts, not bytes -> leave n/l at 0
      reason = DT_OPENCL_FIT_DIMENSION;
    }
    else if(_opencl_get_device_memalloc(devid) < required)
    {
      n = required;
      l = _opencl_get_device_memalloc(devid);
      dt_print(DT_DEBUG_OPENCL,
               "[opencl] trying to allocate %" PRIu64 " MiB of memory while the vRAM has %" PRIu64
               " MiB total\n",
               (uint64_t)(n / (1024 * 1024)), (uint64_t)(l / (1024 * 1024)));
      reason = DT_OPENCL_FIT_ALLOC_LIMIT;
    }
    else if(dt_opencl_get_device_available(devid) < total)
    {
      n = total;
      l = dt_opencl_get_device_available(devid);
      dt_print(DT_DEBUG_OPENCL,
               "[opencl] trying to allocate %" PRIu64 " MiB of memory while the vRAM has %" PRIu64
               " MiB left\n",
               (uint64_t)(n / (1024 * 1024)), (uint64_t)(l / (1024 * 1024)));
      reason = DT_OPENCL_FIT_AVAILABLE;
    }
  }
 
  if(needed) *needed = n;
  if(limit) *limit = l;
  return reason;
}
 
gboolean dt_opencl_image_fits_device(const int devid, const size_t width, const size_t height, const unsigned bpp,
                                const float factor, const size_t overhead)
{
  return dt_opencl_image_fits_device_reason(devid, width, height, bpp, factor, overhead, NULL, NULL)
         == DT_OPENCL_FIT_OK;
}
 
int dt_opencl_dev_roundup_width(int size, const int devid)
{
  const int roundup = darktable.opencl->dev[devid].clroundup_wd;
  return (size % roundup == 0 ? size : (size / roundup + 1) * roundup);
}
int dt_opencl_dev_roundup_height(int size, const int devid)
{
  const int roundup = darktable.opencl->dev[devid].clroundup_ht;
  return (size % roundup == 0 ? size : (size / roundup + 1) * roundup);
}
 
int dt_opencl_is_inited(void)
{
  return darktable.opencl->inited;
}
 
 
int dt_opencl_is_enabled(void)
{
  if(!darktable.opencl->inited) return FALSE;
  return darktable.opencl->enabled;
}
 
 
void dt_opencl_disable(void)
{
  if(!darktable.opencl->inited) return;
  darktable.opencl->enabled = FALSE;
  dt_conf_set_bool("opencl", FALSE);
}
 
 
int dt_opencl_update_settings(void)
{
  dt_opencl_t *cl = darktable.opencl;
  // FIXME: This pulls in prefs every time the pixelpipe runs. Instead have a callback for DT_SIGNAL_PREFERENCES_CHANGE?
  if(!cl->inited) return FALSE;
  const int prefs = dt_conf_get_bool("opencl");
 
  if(cl->enabled != prefs)
  {
    cl->enabled = prefs;
    cl->stopped = 0;
    cl->error_count = 0;
    dt_print(DT_DEBUG_OPENCL, "[opencl_update_enabled] enabled flag set to %s\n", prefs ? "ON" : "OFF");
  }
 
  return (cl->enabled && !cl->stopped);
}
 
 
static void dt_opencl_set_synchronization_timeout(int value)
{
  darktable.opencl->opencl_synchronization_timeout = value;
  dt_print_nts(DT_DEBUG_OPENCL, "[opencl_synchronization_timeout] synchronization timeout set to %d\n", value);
}
 
static void dt_opencl_apply_scheduling_profile()
{
  dt_pthread_mutex_lock(&darktable.opencl->lock);
  dt_opencl_update_priorities();
  dt_opencl_set_synchronization_timeout(dt_conf_get_int("pixelpipe_synchronization_timeout"));
  dt_pthread_mutex_unlock(&darktable.opencl->lock);
}
 
 
cl_event *dt_opencl_events_get_slot(const int devid, const char *tag)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited || devid < 0) return NULL;
  if(!cl->dev[devid].use_events) return NULL;
 
  static const cl_event zeroevent[1]; // implicitly initialized to zero
  cl_event **eventlist = &(cl->dev[devid].eventlist);
  dt_opencl_eventtag_t **eventtags = &(cl->dev[devid].eventtags);
  int *numevents = &(cl->dev[devid].numevents);
  int *maxevents = &(cl->dev[devid].maxevents);
  int *eventsconsolidated = &(cl->dev[devid].eventsconsolidated);
  int *lostevents = &(cl->dev[devid].lostevents);
  int *totalevents = &(cl->dev[devid].totalevents);
  int *totallost = &(cl->dev[devid].totallost);
  int *maxeventslot = &(cl->dev[devid].maxeventslot);
  // if first time called: allocate initial buffers
  if(IS_NULL_PTR(*eventlist))
  {
    int newevents = DT_OPENCL_EVENTLISTSIZE;
    *eventlist = calloc(newevents, sizeof(cl_event));
    *eventtags = calloc(newevents, sizeof(dt_opencl_eventtag_t));
    if(!*eventlist || !*eventtags)
    {
      dt_free(*eventlist);
      dt_free(*eventtags);
      *eventlist = NULL;
      *eventtags = NULL;
      dt_print(DT_DEBUG_OPENCL, "[dt_opencl_events_get_slot] NO eventlist for device %i\n", devid);
      return NULL;
    }
    *maxevents = newevents;
  }
 
  // check if currently highest event slot was actually consumed. If not use it again
  if(*numevents > 0 && !memcmp((*eventlist) + *numevents - 1, zeroevent, sizeof(cl_event)))
  {
    (*lostevents)++;
    (*totallost)++;
    if(!IS_NULL_PTR(tag))
    {
      g_strlcpy((*eventtags)[*numevents - 1].tag, tag, DT_OPENCL_EVENTNAMELENGTH);
    }
    else
    {
      (*eventtags)[*numevents - 1].tag[0] = '\0';
    }
 
    (*totalevents)++;
    return (*eventlist) + *numevents - 1;
  }
 
  // check if we would exceed the number of available event handles. In that case first flush existing handles
  if((*numevents - *eventsconsolidated + 1 > cl->dev[devid].event_handles) || (*numevents == *maxevents))
    (void)dt_opencl_events_flush(devid, 0);
 
  // if no more space left in eventlist: grow buffer
  if(*numevents == *maxevents)
  {
    int newevents = *maxevents + DT_OPENCL_EVENTLISTSIZE;
    cl_event *neweventlist = calloc(newevents, sizeof(cl_event));
    dt_opencl_eventtag_t *neweventtags = calloc(newevents, sizeof(dt_opencl_eventtag_t));
    if(!neweventlist || IS_NULL_PTR(neweventtags))
    {
      dt_print(DT_DEBUG_OPENCL, "[dt_opencl_events_get_slot] NO new eventlist with size %i for device %i\n",
         newevents, devid);
      dt_free(neweventlist);
      dt_free(neweventtags);
      return NULL;
    }
    memcpy(neweventlist, *eventlist, sizeof(cl_event) * *maxevents);
    memcpy(neweventtags, *eventtags, sizeof(dt_opencl_eventtag_t) * *maxevents);
    dt_free(*eventlist);
    dt_free(*eventtags);
    *eventlist = neweventlist;
    *eventtags = neweventtags;
    *maxevents = newevents;
  }
 
  // init next event slot and return it
  (*numevents)++;
  memcpy((*eventlist) + *numevents - 1, zeroevent, sizeof(cl_event));
  if(!IS_NULL_PTR(tag))
  {
    g_strlcpy((*eventtags)[*numevents - 1].tag, tag, DT_OPENCL_EVENTNAMELENGTH);
  }
  else
  {
    (*eventtags)[*numevents - 1].tag[0] = '\0';
  }
 
  (*totalevents)++;
  *maxeventslot = MAX(*maxeventslot, *numevents - 1);
  return (*eventlist) + *numevents - 1;
}
 
 
void dt_opencl_events_reset(const int devid)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited || devid < 0) return;
  if(!cl->dev[devid].use_events) return;
 
  cl_event **eventlist = &(cl->dev[devid].eventlist);
  dt_opencl_eventtag_t **eventtags = &(cl->dev[devid].eventtags);
  int *numevents = &(cl->dev[devid].numevents);
  int *maxevents = &(cl->dev[devid].maxevents);
  int *eventsconsolidated = &(cl->dev[devid].eventsconsolidated);
  int *lostevents = &(cl->dev[devid].lostevents);
  cl_int *summary = &(cl->dev[devid].summary);
 
  if(IS_NULL_PTR(*eventlist) || *numevents == 0) return; // nothing to do
 
  static const cl_event zeroevent[1]; // implicitly initialized to zero
 
  // release all remaining events in eventlist, not to waste resources.
  // Skip NULL handles left behind by failed enqueues: releasing them is
  // pointless and crashes some drivers.
  for(int k = *eventsconsolidated; k < *numevents; k++)
  {
    if(memcmp((*eventlist) + k, zeroevent, sizeof(cl_event)))
      (cl->dlocl->symbols->dt_clReleaseEvent)((*eventlist)[k]);
  }
 
  memset(*eventtags, 0, sizeof(dt_opencl_eventtag_t) * *maxevents);
  *numevents = 0;
  *eventsconsolidated = 0;
  *lostevents = 0;
  *summary = CL_COMPLETE;
  return;
}
 
 
void dt_opencl_events_wait_for(const int devid)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited || devid < 0) return;
  if(!cl->dev[devid].use_events) return;
 
  static const cl_event zeroevent[1]; // implicitly initialized to zero
  cl_event **eventlist = &(cl->dev[devid].eventlist);
  int *numevents = &(cl->dev[devid].numevents);
  int *lostevents = &(cl->dev[devid].lostevents);
  int *totallost = &(cl->dev[devid].totallost);
  int *eventsconsolidated = &(cl->dev[devid].eventsconsolidated);
 
  if(IS_NULL_PTR(*eventlist) || *numevents == 0) return; // nothing to do
 
  // check if last event slot was actually used and correct numevents if needed
  if(!memcmp((*eventlist) + *numevents - 1, zeroevent, sizeof(cl_event)))
  {
    (*numevents)--;
    (*lostevents)++;
    (*totallost)++;
  }
 
  if(*numevents == *eventsconsolidated) return; // nothing to do
 
  assert(*numevents > *eventsconsolidated);
 
  // Wait for all remaining events to terminate, skipping NULL handles. A reserved
  // slot can hold a NULL handle when its enqueue failed (the event is never
  // created, but get_slot had already counted the slot). Passing NULL to
  // clWaitForEvents crashes some drivers and, on those that merely return an
  // error, would abort a batched wait and leave the valid events that follow
  // un-waited-for before flush checks/releases them. We wait on each handle
  // individually rather than gathering them into a temporary array: a heap
  // allocation could fail under the very memory pressure that creates these NULL
  // slots, and returning early on that failure would skip the wait entirely.
  // Risk: might never return in case of OpenCL blocks or endless loops
  // TODO: run clWaitForEvents in separate thread and implement watchdog timer
  for(int k = *eventsconsolidated; k < *numevents; k++)
  {
    if(!memcmp((*eventlist) + k, zeroevent, sizeof(cl_event)))
      continue; // NULL handle from a failed enqueue
 
    cl_int err = (cl->dlocl->symbols->dt_clWaitForEvents)(1, &((*eventlist)[k]));
    if((err != CL_SUCCESS) && (err != CL_INVALID_VALUE))
      dt_vprint(DT_DEBUG_OPENCL, "[dt_opencl_events_wait_for] reported %i for device %i\n",
         err, devid);
  }
}
 
 
cl_int dt_opencl_events_flush(const int devid, const int reset)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited || devid < 0) return FALSE;
  if(!cl->dev[devid].use_events) return FALSE;
 
  cl_event **eventlist = &(cl->dev[devid].eventlist);
  dt_opencl_eventtag_t **eventtags = &(cl->dev[devid].eventtags);
  int *numevents = &(cl->dev[devid].numevents);
  int *eventsconsolidated = &(cl->dev[devid].eventsconsolidated);
  int *lostevents = &(cl->dev[devid].lostevents);
  int *totalsuccess = &(cl->dev[devid].totalsuccess);
 
  cl_int *summary = &(cl->dev[devid].summary);
 
  static const cl_event zeroevent[1]; // implicitly initialized to zero
 
  if(IS_NULL_PTR(*eventlist) || *numevents == 0) return CL_COMPLETE; // nothing to do, no news is good news
 
  // Wait for command queue to terminate (side effect: might adjust *numevents)
  dt_opencl_events_wait_for(devid);
 
  // now check return status and profiling data of all newly terminated events
  for(int k = *eventsconsolidated; k < *numevents; k++)
  {
    // A reserved slot can still hold a NULL handle when the matching enqueue
    // failed: OpenCL does not create the event in that case, but get_slot had
    // already counted the slot. Passing such a NULL handle to the driver
    // crashes some implementations (e.g. NVIDIA on Windows), so treat it as a
    // lost event and skip it.
    if(!memcmp((*eventlist) + k, zeroevent, sizeof(cl_event)))
    {
      (*lostevents)++;
      (*eventsconsolidated)++;
      continue;
    }
 
    cl_int err;
    char *tag = (*eventtags)[k].tag;
    cl_int *retval = &((*eventtags)[k].retval);
 
    // get return value of event
    err = (cl->dlocl->symbols->dt_clGetEventInfo)((*eventlist)[k], CL_EVENT_COMMAND_EXECUTION_STATUS,
                                                  sizeof(cl_int), retval, NULL);
    if(err != CL_SUCCESS)
    {
      dt_print(DT_DEBUG_OPENCL, "[opencl_events_flush] could not get event info for '%s': %i\n",
               tag[0] == '\0' ? "<?>" : tag, err);
    }
    else if(*retval != CL_COMPLETE)
    {
      dt_print(DT_DEBUG_OPENCL, "[opencl_events_flush] execution of '%s' %s: %d\n",
               tag[0] == '\0' ? "<?>" : tag, *retval == CL_COMPLETE ? "was successful" : "failed", *retval);
      *summary = *retval;
    }
    else
      (*totalsuccess)++;
 
    if(darktable.unmuted & DT_DEBUG_PERF)
    {
      // get profiling info of event (only if darktable was called with '-d perf')
      cl_ulong start;
      cl_ulong end;
      cl_int errs = (cl->dlocl->symbols->dt_clGetEventProfilingInfo)(
          (*eventlist)[k], CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
      cl_int erre = (cl->dlocl->symbols->dt_clGetEventProfilingInfo)((*eventlist)[k], CL_PROFILING_COMMAND_END,
                                                                   sizeof(cl_ulong), &end, NULL);
      if(errs == CL_SUCCESS && erre == CL_SUCCESS)
      {
        (*eventtags)[k].timelapsed = end - start;
      }
      else
      {
        (*eventtags)[k].timelapsed = 0;
        (*lostevents)++;
      }
    }
    else
      (*eventtags)[k].timelapsed = 0;
 
    // finally release event to be re-used by driver
    (cl->dlocl->symbols->dt_clReleaseEvent)((*eventlist)[k]);
    (*eventsconsolidated)++;
  }
 
  cl_int result = *summary;
 
  // do we want to get rid of all stored info?
  if(reset)
  {
    // output profiling info if wanted
    if(darktable.unmuted & DT_DEBUG_PERF) dt_opencl_events_profiling(devid, 1);
 
    // reset eventlist structures to empty state
    dt_opencl_events_reset(devid);
  }
 
  return result == CL_COMPLETE ? 0 : result;
}
 
 
void dt_opencl_events_profiling(const int devid, const int aggregated)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited || devid < 0) return;
  if(!cl->dev[devid].use_events) return;
 
  cl_event **eventlist = &(cl->dev[devid].eventlist);
  dt_opencl_eventtag_t **eventtags = &(cl->dev[devid].eventtags);
  int *numevents = &(cl->dev[devid].numevents);
  int *eventsconsolidated = &(cl->dev[devid].eventsconsolidated);
  int *lostevents = &(cl->dev[devid].lostevents);
 
  if(IS_NULL_PTR(*eventlist) || *numevents == 0 || IS_NULL_PTR(*eventtags) || *eventsconsolidated == 0)
    return; // nothing to do
 
  char **tags = malloc(sizeof(char *) * (*eventsconsolidated + 1));
  float *timings = malloc(sizeof(float) * (*eventsconsolidated + 1));
  int items = 1;
  tags[0] = "";
  timings[0] = 0.0f;
 
  // get profiling info and arrange it
  for(int k = 0; k < *eventsconsolidated; k++)
  {
    // if aggregated is TRUE, try to sum up timings for multiple runs of each kernel
    if(aggregated)
    {
      // linear search: this is not efficient at all but acceptable given the limited
      // number of events (ca. 10 - 20)
      int tagfound = -1;
      for(int i = 0; i < items; i++)
      {
        if(!strncmp(tags[i], (*eventtags)[k].tag, DT_OPENCL_EVENTNAMELENGTH))
        {
          tagfound = i;
          break;
        }
      }
 
      if(tagfound >= 0) // tag was already detected before
      {
        // sum up timings
        timings[tagfound] += (*eventtags)[k].timelapsed * 1e-9;
      }
      else // tag is new
      {
        // make new entry
        items++;
        tags[items - 1] = (*eventtags)[k].tag;
        timings[items - 1] = (*eventtags)[k].timelapsed * 1e-9;
      }
    }
 
    else // no aggregated info wanted -> arrange event by event
    {
      items++;
      tags[items - 1] = (*eventtags)[k].tag;
      timings[items - 1] = (*eventtags)[k].timelapsed * 1e-9;
    }
  }
 
  // now display profiling info
  dt_print(DT_DEBUG_OPENCL,
           "[opencl_profiling] profiling device %d ('%s'):\n", devid, cl->dev[devid].name);
 
  float total = 0.0f;
  for(int i = 1; i < items; i++)
  {
    dt_print(DT_DEBUG_OPENCL, "[opencl_profiling] spent %7.4f seconds in %s\n", (double)timings[i],
             tags[i][0] == '\0' ? "<?>" : tags[i]);
    total += timings[i];
  }
  // aggregated timing info for items without tag (if any)
  if(timings[0] != 0.0f)
  {
    dt_print(DT_DEBUG_OPENCL, "[opencl_profiling] spent %7.4f seconds (unallocated)\n", (double)timings[0]);
    total += timings[0];
  }
 
  dt_print(DT_DEBUG_OPENCL,
           "[opencl_profiling] spent %7.4f seconds totally in command queue (with %d event%s missing)\n",
           (double)total, *lostevents, *lostevents == 1 ? "" : "s");
 
  dt_free(timings);
  dt_free(tags);
 
  return;
}
 
static int nextpow2(int n)
{
  int k = 1;
  while (k < n)
    k <<= 1;
  return k;
}
 
// utility function to calculate optimal work group dimensions for a given kernel
// taking device specific restrictions and local memory limitations into account
int dt_opencl_local_buffer_opt(const int devid, const int kernel, dt_opencl_local_buffer_t *factors)
{
  dt_opencl_t *cl = darktable.opencl;
  if(!cl->inited || devid < 0) return FALSE;
 
  size_t maxsizes[3] = { 0 };     // the maximum dimensions for a work group
  size_t workgroupsize = 0;       // the maximum number of items in a work group
  unsigned long localmemsize = 0; // the maximum amount of local memory we can use
  size_t kernelworkgroupsize = 0; // the maximum amount of items in work group for this kernel
 
  int *blocksizex = &factors->sizex;
  int *blocksizey = &factors->sizey;
 
  // initial values must be supplied in sizex and sizey.
  // we make sure that these are a power of 2 and lie within reasonable limits.
  *blocksizex = CLAMP(nextpow2(*blocksizex), 1, 1 << 16);
  *blocksizey = CLAMP(nextpow2(*blocksizey), 1, 1 << 16);
 
  if(dt_opencl_get_work_group_limits(devid, maxsizes, &workgroupsize, &localmemsize) == CL_SUCCESS
     && dt_opencl_get_kernel_work_group_size(devid, kernel, &kernelworkgroupsize) == CL_SUCCESS)
  {
    while(maxsizes[0] < *blocksizex || maxsizes[1] < *blocksizey
       || localmemsize < ((factors->xfactor * (*blocksizex) + factors->xoffset) *
                          (factors->yfactor * (*blocksizey) + factors->yoffset)) * factors->cellsize + factors->overhead
       || workgroupsize < (size_t)(*blocksizex) * (*blocksizey) || kernelworkgroupsize < (size_t)(*blocksizex) * (*blocksizey))
    {
      if(*blocksizex == 1 && *blocksizey == 1) return FALSE;
 
      if(*blocksizex > *blocksizey)
        *blocksizex >>= 1;
      else
        *blocksizey >>= 1;
    }
  }
  else
  {
    dt_print(DT_DEBUG_OPENCL, "[dt_opencl_local_buffer_opt] can not identify resource limits for device %d\n", devid);
    return FALSE;
  }
 
  return TRUE;
}
 
 
#endif
 
// clang-format off
// modelines: These editor modelines have been set for all relevant files by tools/update_modelines.py
// vim: shiftwidth=2 expandtab tabstop=2 cindent
// kate: tab-indents: off; indent-width 2; replace-tabs on; indent-mode cstyle; remove-trailing-spaces modified;
// clang-format on