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-2026 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/paths.h"   // DT_PATH_MAX. conditional-ok: this whole file is inside #ifdef HAVE_OPENCL
#include "common/startup_progress.h"
#include "common/utility.h"   // dt_util_str_replace, used under __APPLE__ only
#include "system/capabilities.h"
#include "system/sys_resources.h"
#include "pixel/bilateralcl.h"
#include "darktable.h"
#include "common/dlopencl.h"
#include "pixel/dwt.h"
#include "common/file_location.h"
#include "pixel/gaussian.h"
#include "pixel/guided_filter.h"
#include "pixel/heal.h"
#include "pixel/interpolation.h"
#include "pixel/locallaplaciancl.h"
#include "system/nvidia_gpus.h"
#include <fcntl.h>   // O_WRONLY/O_CREAT/O_APPEND. conditional-ok: this whole file is inside #ifdef HAVE_OPENCL
#include "system/opencl_drivers_blacklist.h"
#include "common/conf.h"
#include "develop/blend.h"
#include "develop/pixelpipe.h"
#include "caches/pixelpipe_cache.h"
#include "common/times.h"   // dt_get_wtime(). conditional-ok: this whole file is inside #ifdef HAVE_OPENCL
 
#include <assert.h>
#include <locale.h>
#include <stdio.h>
#include <string.h>
#include <strings.h>
 
#include <ctype.h>
#include <errno.h>
#include <glib/gstdio.h>   // g_fopen, g_rename, g_unlink. conditional-ok: this whole file is inside #ifdef HAVE_OPENCL
#include <sys/stat.h>
#include <zlib.h>
 
typedef struct dt_opencl_t
{
  dt_pthread_mutex_t lock;
  int inited;
  int print_statistics;
  int enabled;
  int stopped;
  int num_devs;
  int num_detected_devs;
  int error_count;
  int opencl_synchronization_timeout;
  uint32_t crc;
  int mandatory[5];
  int *dev_priority_image;
  int *dev_priority_preview;
  int *dev_priority_export;
  int *dev_priority_thumbnail;
  dt_opencl_device_t *dev;
  dt_opencl_detected_device_t *detected_devs;
  dt_dlopencl_t *dlocl;
 
  /* NOTE: the per-subsystem OpenCL kernel bundles (blend, bilateral, gaussian, interpolation,
   * local laplacian, dwt, heal, colorspaces, guided filter) used to live here. Each was built
   * by its own subsystem, parked on this struct, and read back from it by that same
   * subsystem -- a round trip that made nine modules' state look like the OpenCL module's.
   * They are file-statics in their owners now; this module only orders their init and free
   * against device setup. */
 
  // Maps every live cl_mem we allocated to the (devid, byte size) we requested,
  // so memory accounting never has to query the driver (clGetMemObjectInfo) about
  // a freshly-created object -- some drivers fault on that under vRAM pressure.
  GHashTable *mem_sizes;
  dt_pthread_mutex_t mem_sizes_lock;
} dt_opencl_t;
 
/* The OpenCL module's state, owned HERE. It used to hang off the application god-struct as
 * `darktable.opencl`, which put every device handle, every kernel bundle and this module's
 * lock one dereference away from any translation unit that included darktable.h. Nothing
 * outside src/common/opencl.* names it now; callers ask questions instead
 * (dt_opencl_get_num_devices(), dt_opencl_get_device_name(), ...) and reserve devices through
 * dt_opencl_reserve_device_*(). */
static dt_opencl_t *_opencl = NULL;
 
/* Internal: neither is called from outside this file. */
static void dt_opencl_cleanup_device(dt_opencl_t *cl, int i);
 
static inline void _opencl_splash_update_compile(const char *programname)
{
  if(IS_NULL_PTR(programname)) return;
 
  // Whether this becomes a splash screen -- and whether one needs opening first -- is the
  // display's business. No "is there a GUI?" test needed: unhandled reports are no-ops.
  dt_startup_progress_report(_("Building OpenCL kernels: %s"), programname);
}
 
static const char *dt_opencl_get_vendor_by_id(unsigned int id);
static const char *_gpu_runtime_of_platform(const char *platform_name);
static int _gpu_runtime_crash_streak(const char *runtime);
static void _gpu_runtime_set_crash_streak(const char *runtime, int value);
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);
 
 
static 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 = _opencl;
  return (!cl->inited || devid < 0) ? 0 : cl->dev[devid].avoid_atomics;
}
 
int dt_opencl_micro_nap(const int devid)
{
  dt_opencl_t *cl = _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 = _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 = _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 = _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);
 
  /* PINNED HOST MEMORY IS THE DEFAULT, and a stored OFF from before this was decided must not
   * keep winning silently.
   *
   * Every buffer that has to cross back to the CPU -- and on a darkroom frame that is at least
   * the displayed output, which Cairo can only read from host memory -- is copied by
   * `dt_opencl_read_host_from_device()`, a BLOCKING read. Into ordinary pageable memory the
   * driver cannot DMA out of the pages it was handed: it stages through an internal pinned
   * buffer and copies twice. Page-locked memory removes the staging copy, and
   * `dt_dev_pixelpipe_cache_sync_cl_buffer()` can additionally take its map/unmap path instead
   * of copying at all, since `dt_opencl_is_pinned_memory()` tests exactly these flags.
   *
   * Measured on a Quadro M2200 over a painting stroke, the same 11.7 MB display buffer, once
   * per frame: 12.43 ms at 0.94 GB/s with this off, 2.04 ms at 5.58 GB/s with it on. It was
   * 38% of the frame and became 7%. The default in `_opencl_device_init()` has always been
   * DT_OPENCL_PINNING_ON; what reached users was a `0` inherited from configurations written
   * before that, which `dt_opencl_read_device_config()` then loaded over it every start.
   *
   * So flip a stored OFF once, per device, and record that we did. The marker is a plain conf
   * key rather than a confgen one on purpose: confgen keys report as existing from the first
   * run, so they cannot answer "has this ever been written", whereas this one is absent until
   * we write it. A user who genuinely wants pinning off sets it again afterwards and keeps it
   * -- the flip happens once, not on every start. A device seen for the first time gets the
   * marker with the default already on, so nothing is flipped for it later.
   *
   * Pinning is not free: page-locked pages cannot be swapped, so the cache flushes them under
   * memory pressure (`flushing cached pinned buffers and retrying`). That is a bounded,
   * handled cost against a copy paid on every frame of every interaction. */
  gchar pin_default_key[256] = { 0 };
  g_snprintf(pin_default_key, sizeof(pin_default_key), "%s/pinned_memory_default_applied", key_device);
  if(!dt_conf_key_exists(pin_default_key))
  {
    if(cl->dev[devid].pinned_memory == DT_OPENCL_PINNING_OFF)
    {
      cl->dev[devid].pinned_memory = DT_OPENCL_PINNING_ON;
      dt_print(DT_DEBUG_OPENCL,
               "[dt_opencl_read_device_config] '%s' had pinned memory transfer off from an older "
               "configuration; enabling it, which is the default\n", cl->dev[devid].cname);
    }
    dt_conf_set_bool(pin_default_key, TRUE);
  }
  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 = _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 = _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 = _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 = _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 = _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].runtime_id = 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(DT_PATH_MAX, sizeof(char));
  char *cachedir = calloc(DT_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[DT_PATH_MAX] = { 0 };
  char *filename = calloc(DT_PATH_MAX, sizeof(char));
  char *confentry = calloc(DT_PATH_MAX, sizeof(char));
  char *binname = calloc(DT_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));
  cl->dev[dev].vendor_id = vendor_id;
 
  /* Which OpenCL runtime publishes this device. A string literal owned by the runtime table,
   * so it outlives the device and is never freed; NULL for a platform we ship no rule for,
   * which simply means no crash can ever be attributed to it. */
  cl->dev[dev].runtime_id = _gpu_runtime_of_platform(platform_name);
  dt_print_nts(DT_DEBUG_OPENCL, "   OPENCL RUNTIME:           %s\n",
               cl->dev[dev].runtime_id ? cl->dev[dev].runtime_id : "unrecognised");
 
  /* Skip this device if ITS runtime is what keeps crashing.
   *
   * Per runtime, not per vendor: Mesa's Rusticl drives AMD and Intel cards alike, so keying on
   * the hardware vendor could neither disable both nor tell a Rusticl panic apart from a fault
   * in AMD's own amdocl on the same card. A machine with an Intel iGPU and an NVIDIA dGPU
   * still keeps the NVIDIA when Intel's compiler is the one faulting -- those are different
   * runtimes, which is exactly what the streak counts. A device whose runtime was never on a
   * crashing stack is untouched however often something else fails.
   *
   * The device name and driver version go in the message because that is what a user needs to
   * act on -- which card, which driver to update. */
  const int runtime_crash_streak = _gpu_runtime_crash_streak(cl->dev[dev].runtime_id);
  if(runtime_crash_streak >= 2)
  {
    dt_print(DT_DEBUG_ALWAYS,
             "[opencl_init] disabling device %d `%s' (%s runtime, driver %s): the last %d"
             " crashes happened inside this OpenCL runtime. Re-enable it in preferences once"
             " the driver is updated; devices on other runtimes are unaffected.\n",
             dev, infostr ? infostr : "?", cl->dev[dev].runtime_id,
             driverversion ? driverversion : "?", runtime_crash_streak);
    cl->dev[dev].disabled |= 1;
    res = -1;
    goto end;
  }
 
  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, (dt_get_debug_flags() & 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, DT_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, DT_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;
 
  /* Every header a .cl may include, so a change to one invalidates the cached device
   * binaries that were built from it -- an omission here does not fail, it silently keeps
   * running yesterday's kernel. lensserious_eval.h is LensSerious's evaluator, installed
   * beside the kernels by data/kernels/CMakeLists.txt. */
  const char *clincludes[DT_OPENCL_MAX_INCLUDES] = { "rgb_norms.h", "noise_generator.h", "color_conversion.h", "colorspaces.cl", "colorspace.h", "common.h", "lensserious_eval.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_flush();
  }
 
  // 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", DT_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;
}
 
/* OpenCL runtimes, by the name their shared object carries in a crash backtrace and the
 * CL_PLATFORM_NAME of the platform they publish.
 *
 * The module name is deliberately the RUNTIME's own object, never the ICD loader (OpenCL.dll /
 * libOpenCL.so): the loader is on the stack of every OpenCL crash, including ones that are
 * entirely ours.
 *
 * The key is the RUNTIME, not the hardware vendor. A runtime is what crashes, and a runtime is
 * what a device can be moved off; the vendor is neither. Mesa's Rusticl is what forces the
 * distinction (issue #1226): it drives AMD radeonsi and Intel iris cards alike, so a
 * vendor-keyed streak could neither cover both nor tell a Rusticl panic apart from a fault in
 * AMD's own amdocl on the very same card. A device belongs to exactly one platform and a
 * platform is published by exactly one runtime, so CL_PLATFORM_NAME is the honest way back
 * from a device to the thing that crashed.
 *
 * Rows may share a conf_suffix: one runtime can ship several objects, and either can be the one
 * on the stack (Intel's compiler and its runtime proper are separate libraries). The suffix,
 * not the row, is the identity; a platform lookup stops at the first row that matches.
 *
 * Platform matching is a case-insensitive substring test and it fails OPEN -- an unrecognised
 * platform gets no streak and is never disabled, which is what every platform did before this
 * table existed. That is the safe direction, but it means a runtime that renames its platform
 * silently stops being covered, so each device logs the runtime it resolved to under
 * `-d opencl` and a report tells us which string to add.
 *
 * The platform strings for rusticl, Intel, NVIDIA and Clover are sampled from clinfo; AMD's is
 * the long-standing name of its APP platform under both the proprietary stack and ROCm, and is
 * the one row here not read off a live machine. Nothing else publishes a platform name
 * containing "AMD", so the short substring costs no precision and buys coverage if AMD ever
 * decorates the rest of the string. */
typedef struct _gpu_runtime_t
{
  const char *module;       // substring to look for in a crash backtrace
  const char *platform;     // substring of the CL_PLATFORM_NAME it publishes
  const char *conf_suffix;  // stable fragment of the per-runtime conf key
} _gpu_runtime_t;
 
static const _gpu_runtime_t _gpu_runtimes[] = {
  { "RusticlOpenCL",    "rusticl",         "rusticl" },  // issue #1226
  { "libMesaOpenCL",    "Clover",          "clover" },   // Rusticl's predecessor, same exposure
  { "amdocl",           "AMD",             "amd" },      // Sentry 129862572, 141634558
  { "libigdfcl",        "Intel(R) OpenCL", "intel" },    // Sentry 129978857
  { "libigdrcl",        "Intel(R) OpenCL", "intel" },
  { "intelocl",         "Intel(R) OpenCL", "intel" },
  { "libnvidia-opencl", "NVIDIA",          "nvidia" },
  { "nvopencl",         "NVIDIA",          "nvidia" },
  { NULL, NULL, NULL }
};
 
/* The runtime publishing `platform_name`, as its conf-key suffix; NULL if we ship no rule for
 * it. Case-insensitive, following dt_opencl_check_driver_blacklist(). */
static const char *_gpu_runtime_of_platform(const char *platform_name)
{
  if(IS_NULL_PTR(platform_name)) return NULL;
 
  gchar *haystack = g_ascii_strdown(platform_name, -1);
  const char *runtime = NULL;
  for(int i = 0; _gpu_runtimes[i].module && IS_NULL_PTR(runtime); i++)
  {
    gchar *needle = g_ascii_strdown(_gpu_runtimes[i].platform, -1);
    if(strstr(haystack, needle)) runtime = _gpu_runtimes[i].conf_suffix;
    g_free(needle);
  }
  g_free(haystack);
  return runtime;
}
 
static int _gpu_runtime_crash_streak(const char *runtime)
{
  if(IS_NULL_PTR(runtime)) return 0;
  char key[128];
  snprintf(key, sizeof(key), "opencl_driver_crash_streak_%s", runtime);
  return dt_conf_get_int(key);
}
 
static void _gpu_runtime_set_crash_streak(const char *runtime, int value)
{
  if(IS_NULL_PTR(runtime)) return;
  char key[128];
  snprintf(key, sizeof(key), "opencl_driver_crash_streak_%s", runtime);
  dt_conf_set_int(key, value);
}
 
/* Path of the driver-crash record. Built at init, because dt_opencl_note_crash_backtrace()
 * runs in a crash handler and cannot allocate; empty until then, and that function no-ops
 * while it is. */
static char _driver_crash_marker[DT_PATH_MAX] = { 0 };
 
static void _opencl_driver_crash_marker_path(char *buf, size_t buf_len)
{
  char cachedir[DT_PATH_MAX] = { 0 };
  dt_loc_get_user_cache_dir(cachedir, sizeof(cachedir));
  char *marker = g_build_filename(cachedir, "gpu-driver-crash", NULL);
  g_strlcpy(buf, marker, buf_len);
  g_free(marker);
}
 
void dt_opencl_note_crash_backtrace(const char *backtrace, size_t backtrace_len)
{
  if(_driver_crash_marker[0] == '\0' || IS_NULL_PTR(backtrace) || backtrace_len == 0) return;
 
  /* Record WHICH runtime, not just that one crashed: the next start disables only the devices
   * that runtime publishes. Writing the module's own compile-time string keeps this handler
   * free of any formatting -- the module is mapped back to its runtime at the next start, off
   * the crash path. */
  const char *culprit = NULL;
  for(int i = 0; _gpu_runtimes[i].module && IS_NULL_PTR(culprit); i++)
    if(strstr(backtrace, _gpu_runtimes[i].module)) culprit = _gpu_runtimes[i].module;
 
  if(IS_NULL_PTR(culprit)) return;
 
  /* open/write/close only: this runs on a dying process whose heap may be corrupt, so no
   * allocation, no locks, no conf. One line per crash. */
  const int fd = g_open(_driver_crash_marker, O_WRONLY | O_CREAT | O_APPEND, 0600);
  if(fd < 0) return;
  const ssize_t w1 = write(fd, culprit, strlen(culprit));
  const ssize_t w2 = write(fd, "\n", 1);
  (void)w1; (void)w2;
  close(fd);
}
 
/* Fold the recorded driver crashes into the per-vendor streaks, consuming the record.
 *
 * Runs early in dt_opencl_init(), where allocation and conf are safe again. Each line names
 * the module that was on the crashing stack; it is mapped back to the runtime that ships it,
 * so the streak that grows is that runtime's alone. Lines written by an older build name the
 * same modules, so a marker left across an upgrade still folds correctly. */
static void _opencl_fold_driver_crashes(void)
{
  if(_driver_crash_marker[0] == '\0') return;
 
  gchar *contents = NULL;
  gsize length = 0;
  if(!g_file_get_contents(_driver_crash_marker, &contents, &length, NULL)) return;
  g_unlink(_driver_crash_marker);
 
  if(!IS_NULL_PTR(contents) && length > 0)
  {
    gchar **lines = g_strsplit(contents, "\n", -1);
    for(int l = 0; lines[l]; l++)
    {
      if(lines[l][0] == '\0') continue;
      for(int i = 0; _gpu_runtimes[i].module; i++)
      {
        if(strcmp(lines[l], _gpu_runtimes[i].module) != 0) continue;
        const char *runtime = _gpu_runtimes[i].conf_suffix;
        _gpu_runtime_set_crash_streak(runtime, _gpu_runtime_crash_streak(runtime) + 1);
        break;
      }
    }
    g_strfreev(lines);
  }
  g_free(contents);
}
 
void dt_opencl_init(const gboolean exclude_opencl, const gboolean print_statistics)
{
  if(_opencl) return;
  _opencl = (dt_opencl_t *)calloc(1, sizeof(dt_opencl_t));
  if(IS_NULL_PTR(_opencl)) return;
  dt_opencl_t *cl = _opencl;
  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;
 
  /* NOTE: this block sits ABOVE the locale save below on purpose. It can return early, and a
   * plain return past that point would skip the finally: block that restores the locale,
   * leaving the whole session in "C" -- wrong number and date formatting everywhere, for
   * exactly the users already suffering driver crashes. Jumping to finally: instead is not an
   * option either: that label frees all_platforms, platform_name and friends, which are
   * declared further down, so a goto from here would jump over their initialisers and free
   * indeterminate pointers. Returning before anything needs unwinding avoids both. */
  /* Fold in any crash that happened inside a GPU driver last time.
   *
   * A vendor OpenCL runtime that faults on its own thread is not something this code can fix
   * -- the stack holds no frame of ours (Sentry 129862572: 169 crashes from 10 users entirely
   * inside amdocl64.dll; 129978857: 37 from 17 users, Intel's compiler longjmp-ing out of a
   * signal handler into glibc's fortify check). What it CAN do is stop walking into the same
   * driver every launch: ~17 and ~2 crashes per user means the same wall every start.
   *
   * The decision is per DEVICE, further down, once each device's vendor is known -- a machine
   * with an Intel iGPU and an NVIDIA dGPU must not lose the NVIDIA because Intel's compiler
   * faulted. Nothing is disabled here. */
  _opencl_driver_crash_marker_path(_driver_crash_marker, sizeof(_driver_crash_marker));
  _opencl_fold_driver_crashes();
 
  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;
 
  // Before the locale save: only finally: restores it.
  if(exclude_opencl) return;
 
  // 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_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");
    dt_develop_blend_init_cl_global();
    dt_bilateral_init_cl_global();
    dt_gaussian_init_cl_global();
    dt_interpolation_init_cl_global();
    dt_local_laplacian_init_cl_global();
    dt_dwt_init_cl_global();
    dt_heal_init_cl_global();
    dt_colorspaces_init_cl_global();
    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;
}
 
static 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 && (dt_get_debug_flags() & 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_clear_driver_crash_streak(void)
{
  /* Clear the driver-crash streak only if OpenCL actually RAN and we still reached a clean
   * shutdown: that is the one thing showing the driver behaving. Surviving with OpenCL
   * switched off proves nothing about it, and clearing on that would re-enable it next launch
   * and crash again -- halving the crash rate instead of stopping it.
   *
   * This touches conf, so it must run while conf is still alive -- i.e. BEFORE
   * dt_conf_cleanup(), which is why it is not part of dt_opencl_cleanup(): that one runs after
   * conf has been saved and freed, where a conf access is both a use-after-free and a write
   * that can never be persisted. */
  dt_opencl_t *cl = _opencl;
  if(IS_NULL_PTR(cl) || !cl->inited || !cl->enabled || IS_NULL_PTR(cl->dev)) return;
 
  /* The counters are per runtime (the same keys _gpu_runtime_crash_streak() gates a device
   * on), and only a runtime whose device RAN is cleared: a device this session skipped because
   * of its own streak has not shown anything. */
  for(int i = 0; i < cl->num_devs; i++)
  {
    if(cl->dev[i].disabled) continue;
    const char *runtime = cl->dev[i].runtime_id;
    if(_gpu_runtime_crash_streak(runtime) != 0)
      _gpu_runtime_set_crash_streak(runtime, 0);
  }
}
 
void dt_opencl_cleanup(void)
{
  dt_opencl_t *cl = _opencl;
  if(IS_NULL_PTR(cl)) return;
  if(cl->inited)
  {
    dt_develop_blend_free_cl_global();
    dt_bilateral_free_cl_global();
    dt_gaussian_free_cl_global();
    dt_interpolation_free_cl_global();
    dt_local_laplacian_free_cl_global();
    dt_dwt_free_cl_global();
    dt_heal_free_cl_global();
    dt_colorspaces_free_cl_global();
    dt_guided_filter_free_cl_global();
 
    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);
 
  dt_free(_opencl);
  _opencl = NULL;
}
 
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 = _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 = _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 = _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 = _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_reserve_device_for_pipe(const int pipetype)
{
  dt_opencl_t *cl = _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_release_device(const int devid)
{
  dt_opencl_t *cl = _opencl;
  if(!cl->inited) return;
  if(devid < 0 || devid >= cl->num_devs) return;
  dt_pthread_mutex_BAD_unlock(&cl->dev[devid].lock);
}
 
/* The same per-device lock dt_opencl_reserve_device_for_pipe() takes, for callers that
 * already know which device they need -- typically to touch that device's cl_mem or its
 * eventlist, which races the pipe using it otherwise. */
void dt_opencl_reserve_device_by_id(const int devid)
{
  dt_opencl_t *cl = _opencl;
  if(!cl->inited) return;
  if(devid < 0 || devid >= cl->num_devs) return;
  // paired with dt_opencl_release_device()'s dt_pthread_mutex_BAD_unlock(): the lock is held
  // across this function boundary and released by a different function, which the analyzer
  // can't model -- same reason dt_opencl_try_reserve_device_by_id() already uses the BAD variant.
  dt_pthread_mutex_BAD_lock(&cl->dev[devid].lock);
}
 
int dt_opencl_try_reserve_device_by_id(const int devid)
{
  dt_opencl_t *cl = _opencl;
  if(!cl->inited) return 1;
  if(devid < 0 || devid >= cl->num_devs) return 1;
  return dt_pthread_mutex_BAD_trylock(&cl->dev[devid].lock);
}
 
int dt_opencl_get_num_devices(void)
{
  dt_opencl_t *cl = _opencl;
  if(!cl || !cl->inited) return 0;
  return cl->num_devs;
}
 
const char *dt_opencl_get_device_name(const int devid)
{
  dt_opencl_t *cl = _opencl;
  if(!cl || !cl->inited || IS_NULL_PTR(cl->dev)) return NULL;
  if(devid < 0 || devid >= cl->num_devs) return NULL;
  return cl->dev[devid].name;
}
 
gboolean dt_opencl_get_device_max_image_size(const int devid, int *width, int *height)
{
  dt_opencl_t *cl = _opencl;
  if(!cl || !cl->inited || IS_NULL_PTR(cl->dev)) return FALSE;
  if(devid < 0 || devid >= cl->num_devs) return FALSE;
  if(width) *width = cl->dev[devid].max_image_width;
  if(height) *height = cl->dev[devid].max_image_height;
  return TRUE;
}
 
size_t dt_opencl_get_device_max_global_mem(const int devid)
{
  dt_opencl_t *cl = _opencl;
  if(!cl || !cl->inited || IS_NULL_PTR(cl->dev)) return 0;
  if(devid < 0 || devid >= cl->num_devs) return 0;
  return (size_t)cl->dev[devid].max_global_mem;
}
 
size_t dt_opencl_get_device_memory_in_use(const int devid)
{
  dt_opencl_t *cl = _opencl;
  if(!cl || !cl->inited || IS_NULL_PTR(cl->dev)) return 0;
  if(devid < 0 || devid >= cl->num_devs) return 0;
  return cl->dev[devid].memory_in_use;
}
 
size_t dt_opencl_get_device_headroom(const int devid)
{
  dt_opencl_t *cl = _opencl;
  if(!cl || !cl->inited || IS_NULL_PTR(cl->dev)) return 0;
  if(devid < 0 || devid >= cl->num_devs) return 0;
  // In MiB, as dt_opencl_check_tuning() reads it.
  return cl->dev[devid].forced_headroom * 1024 * 1024;
}
 
int dt_opencl_report_pipe_error(void)
{
  dt_opencl_t *cl = _opencl;
  if(!cl) return 2;
 
  cl->error_count++;
  if(cl->error_count < DT_OPENCL_MAX_ERRORS) return 1;
 
  /* Too many failures: OpenCL is finished for this session. */
  cl->stopped = 1;
  dt_capabilities_remove("opencl");
  return 2;
}
 
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);
    fclose(f); // the open succeeded; only the stat failed, and the handle is still ours
    return NULL;
  }
  return f;
}
 
 
void dt_opencl_md5sum(const char **files, char **md5sums)
{
  char kerneldir[DT_PATH_MAX] = { 0 };
  char filename[DT_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 = _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[DT_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[DT_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);
        if(IS_NULL_PTR(cached_content))
        {
          dt_print(DT_DEBUG_OPENCL,
                   "[opencl_load_program] could not allocate %zu bytes for cached binary '%s'!\n",
                   cached_filesize + 1, binname);
          fclose(cached);
          return 0;
        }
        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[DT_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 = _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[DT_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;
          const size_t bytes_written = fwrite(binaries[i], sizeof(char), binary_sizes[i], f);
          fclose(f);
          if(bytes_written != binary_sizes[i]) goto ret;
 
          /* Name the binary after the program (e.g. basic.cl.bin), pointing at the
           * md5sum-named file we just wrote.
           *
           * This used to chdir() into cachedir so it could name the link by its basename
           * alone, and chdir() back afterwards. Two things were wrong with that. The
           * working directory is process-global state, and this runs on whichever thread
           * is building kernels while the rest of the application is live; and the
           * restoring chdir() sat AFTER two `goto ret' paths, so any failure in between
           * left the whole process parked in the kernel cache directory for the rest of
           * the session. Neither is needed: only the LINK path has to be absolute, while
           * the symlink target stays relative -- which is what the chdir was buying. */
          gchar *bname = g_path_get_basename(binname);
          char link_name[DT_PATH_MAX] = { 0 };
          snprintf(link_name, sizeof(link_name), "%s" G_DIR_SEPARATOR_S "%s", cachedir, bname);
#if defined(_WIN32)
          /* CreateSymbolicLink() needs admin privileges on Windows, which we neither
           * want nor need, so the cached binary is simply renamed into place. */
          char final_name[DT_PATH_MAX] = { 0 };
          snprintf(final_name, sizeof(final_name), "%s.%s", link_name, md5sum);
          const int link_failed = (g_rename(link_dest, final_name) != 0);
#else
          /* Relative target, absolute link path: the link keeps working if the cache
           * directory is ever moved, and we never touch the working directory. */
          g_unlink(link_name);
          const int link_failed = (symlink(md5sum, link_name) != 0);
#endif 
          g_free(bname);
          if(link_failed)
          {
            dt_print(DT_DEBUG_OPENCL, "[opencl_build_program] could not name the cached binary `%s': %s\n",
                     link_name, strerror(errno));
            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 = _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 = _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 = _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 = _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 = _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 = _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 = _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(dt_get_debug_flags() & 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(!(_opencl && _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(!(_opencl && _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(!(_opencl && _opencl->inited)) return -1;
 
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Read Image (from device to host)]");
 
  return (_opencl->dlocl->symbols->dt_clEnqueueReadImage)(_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(!(_opencl && _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(!(_opencl && _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(!(_opencl && _opencl->inited)) return -1;
 
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Write Image (from host to device)]");
 
  return (_opencl->dlocl->symbols->dt_clEnqueueWriteImage)(_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(!(_opencl && _opencl->inited) || devid < 0) return -1;
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Copy Image (on device)]");
  cl_int err = (_opencl->dlocl->symbols->dt_clEnqueueCopyImage)(
      _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(!(_opencl && _opencl->inited)) return -1;
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Copy Image to Buffer (on device)]");
  cl_int err = (_opencl->dlocl->symbols->dt_clEnqueueCopyImageToBuffer)(
      _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(!(_opencl && _opencl->inited)) return -1;
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Copy Buffer to Image (on device)]");
  cl_int err = (_opencl->dlocl->symbols->dt_clEnqueueCopyBufferToImage)(
      _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(!(_opencl && _opencl->inited)) return -1;
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Copy Buffer to Buffer (on device)]");
  cl_int err = (_opencl->dlocl->symbols->dt_clEnqueueCopyBuffer)(_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(!(_opencl && _opencl->inited)) return -1;
 
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Read Buffer (from device to host)]");
 
  return (_opencl->dlocl->symbols->dt_clEnqueueReadBuffer)(
      _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(!(_opencl && _opencl->inited)) return -1;
 
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Write Buffer (from host to device)]");
 
  return (_opencl->dlocl->symbols->dt_clEnqueueWriteBuffer)(
      _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(!(_opencl && _opencl->inited) || devid < 0) return NULL;
  cl_int err;
  cl_mem dev = (_opencl->dlocl->symbols->dt_clCreateBuffer)(
      _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(!(_opencl && _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 = (_opencl->dlocl->symbols->dt_clCreateImage2D)(
      _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(!(_opencl && _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);
 
  (_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(!(_opencl && _opencl->inited)) return NULL;
  cl_int err;
  void *ptr;
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Map Buffer]");
  ptr = (_opencl->dlocl->symbols->dt_clEnqueueMapBuffer)(
      _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(!(_opencl && _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 = (_opencl->dlocl->symbols->dt_clEnqueueMapImage)(
      _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(!(_opencl && _opencl->inited)) return -1;
  cl_event *eventp = dt_opencl_events_get_slot(devid, "[Unmap Mem Object]");
  cl_int err = (_opencl->dlocl->symbols->dt_clEnqueueUnmapMemObject)(
      _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(!(_opencl && _opencl->inited) || devid < 0) return NULL;
  cl_int err;
  cl_mem dev = NULL;
  for(int attempt = 0; attempt < 2; attempt++)
  {
    dev = (_opencl->dlocl->symbols->dt_clCreateImage2D)(_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(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(!(_opencl && _opencl->inited)) return NULL;
  cl_int err;
  cl_mem buf = NULL;
  for(int attempt = 0; attempt < 2; attempt++)
  {
    buf = (_opencl->dlocl->symbols->dt_clCreateBuffer)(_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(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 = (_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 = (_opencl->dlocl->symbols->dt_clGetMemObjectInfo)(mem, CL_MEM_CONTEXT, sizeof(context), &context, NULL);
  if(err != CL_SUCCESS)
    return -1;
 
  for(int devid = 0; devid < _opencl->num_devs; devid++)
  {
    if(_opencl->dev[devid].context == context)
      return devid;
  }
 
  return -1;
}
 
cl_mem_flags dt_opencl_get_mem_flags(cl_mem mem)
{
  if(!(_opencl && _opencl->inited) || IS_NULL_PTR(mem)) return 0;
  cl_mem_flags flags = 0;
  cl_int err = (_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 = (_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 = (_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 = (_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(&_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(_opencl->mem_sizes, mem, rec);
    dt_pthread_mutex_unlock(&_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(&_opencl->mem_sizes_lock);
    dt_opencl_mem_record_t *rec = (dt_opencl_mem_record_t *)g_hash_table_lookup(_opencl->mem_sizes, mem);
    if(rec)
    {
      devid = rec->devid;
      size = rec->size;
      g_hash_table_remove(_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(&_opencl->mem_sizes_lock);
  }
 
  if(devid < 0)
    return;
 
  if(action == OPENCL_MEMORY_ADD)
    _opencl->dev[devid].memory_in_use += size;
  else
    _opencl->dev[devid].memory_in_use =
      (_opencl->dev[devid].memory_in_use > size)
        ? (_opencl->dev[devid].memory_in_use - size)
        : 0;
 
  _opencl->dev[devid].peak_memory = MAX(_opencl->dev[devid].peak_memory,
                                                 _opencl->dev[devid].memory_in_use);
 
  if((dt_get_debug_flags() & DT_DEBUG_MEMORY) && (dt_get_debug_flags() & DT_DEBUG_OPENCL))
    dt_print(DT_DEBUG_OPENCL,
              "[opencl memory] device %d: %" G_GSIZE_FORMAT " bytes (%.1f MB) in use\n", devid, _opencl->dev[devid].memory_in_use,
                                      (float)_opencl->dev[devid].memory_in_use/(1024*1024));
}
 
gboolean dt_opencl_get_mem_record(cl_mem mem, int *devid, size_t *size)
{
  if(IS_NULL_PTR(mem) || !(_opencl && _opencl->inited)) return FALSE;
  dt_pthread_mutex_lock(&_opencl->mem_sizes_lock);
  const dt_opencl_mem_record_t *rec = (const dt_opencl_mem_record_t *)g_hash_table_lookup(_opencl->mem_sizes, mem);
  const gboolean found = !IS_NULL_PTR(rec);
  if(found)
  {
    *devid = rec->devid;
    *size = rec->size;
  }
  dt_pthread_mutex_unlock(&_opencl->mem_sizes_lock);
  return found;
}
 
void dt_opencl_check_tuning(const int devid)
{
  dt_opencl_t *cl = _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(!(_opencl && _opencl->inited) || devid < 0) return 0;
  const cl_ulong limit = _opencl->dev[devid].used_available;
  const size_t in_use = _opencl->dev[devid].memory_in_use;
  cl_ulong available = (limit > in_use) ? (limit - in_use) : 0;
 
  /* A host-unified device is bounded by two DIFFERENT things, and they must not be conflated.
   *
   * 1. Free space. Its "vRAM" is system RAM, shared with every other process and with the host
   *    side of our own pipeline, so max_global_mem is the size of the machine rather than a
   *    budget -- a P630 reports 28.75 GiB on a 31 GiB box. dt_get_system_available_mem() already
   *    measures what is actually left (free plus reclaimable, cgroup limits included), so clamp
   *    to it. A 0 return means the platform will not say; that is "no information", not "no
   *    memory", so the device's own accounting stands.
   *
   * 2. How much the driver can have in flight at once, which is NOT a free-space question and
   *    does not get better when RAM is free. Measured on the P630 with 22 GiB genuinely
   *    available: an ~8 GiB working set fails, the pipe loses OpenCL entirely and drops to CPU;
   *    ~4 GiB fails mid-tile; ~2 GiB completes. The ceiling is the GPU aperture / address space
   *    per context, which OpenCL publishes no query for. max_mem_alloc is the closest thing it
   *    does publish -- NEO reports min(global/4, 4 GiB), which tracks that aperture -- so half
   *    of it is used as a working-set ceiling. It is a proxy chosen to match the measurement,
   *    not a measurement of anything, and it is written here as its own clamp so it cannot be
   *    misread as free space.
   *
   * This lives in the reader rather than in dt_opencl_check_tuning() because the tilers and the
   * fit checks call it while deciding: limit 1 is then re-probed at each decision instead of
   * frozen once per pipe run, before any of the allocations it governs. The probe caches for
   * tens of milliseconds, which is what makes that affordable.
   *
   * Discrete devices never enter here: their max_global_mem really is dedicated vRAM. */
  if(_opencl->dev[devid].host_unified_memory)
  {
    const size_t reserved = (size_t)_opencl->dev[devid].forced_headroom * 1024 * 1024;
 
    const size_t system_available = dt_get_system_available_mem();
    if(system_available > 0)
      available = MIN(available, (cl_ulong)((system_available > reserved) ? system_available - reserved : 0));
 
    available = MIN(available, (cl_ulong)(_opencl->dev[devid].max_mem_alloc / 2));
  }
  return available;
}
 
static cl_ulong _opencl_get_device_memalloc(const int devid)
{
  return _opencl->dev[devid].max_mem_alloc;
}
 
cl_ulong dt_opencl_get_device_memalloc(const int devid)
{
  if(!(_opencl && _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 = _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 = _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 = _opencl->dev[devid].clroundup_ht;
  return (size % roundup == 0 ? size : (size / roundup + 1) * roundup);
}
 
int dt_opencl_is_inited(void)
{
  return (_opencl && _opencl->inited);
}
 
 
int dt_opencl_is_enabled(void)
{
  if(!(_opencl && _opencl->inited)) return FALSE;
  return _opencl->enabled;
}
 
 
void dt_opencl_disable(void)
{
  if(!(_opencl && _opencl->inited)) return;
  _opencl->enabled = FALSE;
  dt_conf_set_bool("opencl", FALSE);
}
 
 
int dt_opencl_update_settings(void)
{
  dt_opencl_t *cl = _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)
{
  _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(&_opencl->lock);
  dt_opencl_update_priorities();
  dt_opencl_set_synchronization_timeout(dt_conf_get_int("pixelpipe_synchronization_timeout"));
  dt_pthread_mutex_unlock(&_opencl->lock);
}
 
 
cl_event *dt_opencl_events_get_slot(const int devid, const char *tag)
{
  dt_opencl_t *cl = _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 = _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 = _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 = _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(dt_get_debug_flags() & 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(dt_get_debug_flags() & 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 = _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 = _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