hip.h

// Start of backends/hip.h.

// Forward declarations.
// Invoked by setup_opencl() after the platform and device has been
// found, but before the program is loaded.  Its intended use is to
// tune constants based on the selected platform and device.
static void set_tuning_params(struct futhark_context* ctx);
static char* get_failure_msg(int failure_idx, int64_t args[]);

#define HIP_SUCCEED_FATAL(x) hip_api_succeed_fatal(x, #x, __FILE__, __LINE__)
#define HIP_SUCCEED_NONFATAL(x) hip_api_succeed_nonfatal(x, #x, __FILE__, __LINE__)
#define HIPRTC_SUCCEED_FATAL(x) hiprtc_api_succeed_fatal(x, #x, __FILE__, __LINE__)
#define HIPRTC_SUCCEED_NONFATAL(x) hiprtc_api_succeed_nonfatal(x, #x, __FILE__, __LINE__)
// Take care not to override an existing error.
#define HIP_SUCCEED_OR_RETURN(e) {             \
    char *serror = HIP_SUCCEED_NONFATAL(e);    \
    if (serror) {                               \
      if (!ctx->error) {                        \
        ctx->error = serror;                    \
        return bad;                             \
      } else {                                  \
        free(serror);                           \
      }                                         \
    }                                           \
  }

// HIP_SUCCEED_OR_RETURN returns the value of the variable 'bad' in
// scope.  By default, it will be this one.  Create a local variable
// of some other type if needed.  This is a bit of a hack, but it
// saves effort in the code generator.
static const int bad = 1;

static inline void hip_api_succeed_fatal(hipError_t res, const char *call,
                                         const char *file, int line) {
  if (res != hipSuccess) {
    const char *err_str = hipGetErrorString(res);
    if (err_str == NULL) { err_str = "Unknown"; }
    futhark_panic(-1, "%s:%d: HIP call\n  %s\nfailed with error code %d (%s)\n",
                  file, line, call, res, err_str);
  }
}

static char* hip_api_succeed_nonfatal(hipError_t res, const char *call,
                                      const char *file, int line) {
  if (res != hipSuccess) {
    const char *err_str = hipGetErrorString(res);
    if (err_str == NULL) { err_str = "Unknown"; }
    return msgprintf("%s:%d: HIP call\n  %s\nfailed with error code %d (%s)\n",
                     file, line, call, res, err_str);
  } else {
    return NULL;
  }
}

static inline void hiprtc_api_succeed_fatal(hiprtcResult res, const char *call,
                                           const char *file, int line) {
  if (res != HIPRTC_SUCCESS) {
    const char *err_str = hiprtcGetErrorString(res);
    futhark_panic(-1, "%s:%d: HIPRTC call\n  %s\nfailed with error code %d (%s)\n",
                  file, line, call, res, err_str);
  }
}

static char* hiprtc_api_succeed_nonfatal(hiprtcResult res, const char *call,
                                        const char *file, int line) {
  if (res != HIPRTC_SUCCESS) {
    const char *err_str = hiprtcGetErrorString(res);
    return msgprintf("%s:%d: HIPRTC call\n  %s\nfailed with error code %d (%s)\n",
                     file, line, call, res, err_str);
  } else {
    return NULL;
  }
}

struct futhark_context_config {
  int in_use;
  int debugging;
  int profiling;
  int logging;
  const char *cache_fname;
  int num_tuning_params;
  int64_t *tuning_params;
  const char** tuning_param_names;
  const char** tuning_param_vars;
  const char** tuning_param_classes;
  // Uniform fields above.

  char* program;
  int num_build_opts;
  const char **build_opts;

  const char *preferred_device;
  int preferred_device_num;

  size_t default_block_size;
  size_t default_grid_size;
  size_t default_tile_size;
  size_t default_reg_tile_size;
  size_t default_threshold;

  int default_block_size_changed;
  int default_grid_size_changed;
  int default_tile_size_changed;
};

static void backend_context_config_setup(struct futhark_context_config *cfg) {
  cfg->num_build_opts = 0;
  cfg->build_opts = (const char**) malloc(sizeof(const char*));
  cfg->build_opts[0] = NULL;
  cfg->preferred_device_num = 0;
  cfg->preferred_device = "";
  cfg->program = strconcat(gpu_program);

  cfg->default_block_size = 256;
  cfg->default_grid_size = 0; // Set properly later.
  cfg->default_tile_size = 32;
  cfg->default_reg_tile_size = 2;
  cfg->default_threshold = 32*1024;

  cfg->default_block_size_changed = 0;
  cfg->default_grid_size_changed = 0;
  cfg->default_tile_size_changed = 0;
}

static void backend_context_config_teardown(struct futhark_context_config* cfg) {
  free(cfg->build_opts);
  free(cfg->program);
}

void futhark_context_config_add_build_option(struct futhark_context_config *cfg, const char *opt) {
  cfg->build_opts[cfg->num_build_opts] = opt;
  cfg->num_build_opts++;
  cfg->build_opts = (const char **) realloc(cfg->build_opts, (cfg->num_build_opts + 1) * sizeof(const char *));
  cfg->build_opts[cfg->num_build_opts] = NULL;
}

void futhark_context_config_set_device(struct futhark_context_config *cfg, const char *s) {
  int x = 0;
  if (*s == '#') {
    s++;
    while (isdigit(*s)) {
      x = x * 10 + (*s++)-'0';
    }
    // Skip trailing spaces.
    while (isspace(*s)) {
      s++;
    }
  }
  cfg->preferred_device = s;
  cfg->preferred_device_num = x;
}


const char* futhark_context_config_get_program(struct futhark_context_config *cfg) {
  return cfg->program;
}

void futhark_context_config_set_program(struct futhark_context_config *cfg, const char *s) {
  cfg->program = strdup(s);
}

void futhark_context_config_set_default_group_size(struct futhark_context_config *cfg, int size) {
  cfg->default_block_size = size;
  cfg->default_block_size_changed = 1;
}

void futhark_context_config_set_default_num_groups(struct futhark_context_config *cfg, int num) {
  cfg->default_grid_size = num;
  cfg->default_grid_size_changed = 1;
}

void futhark_context_config_set_default_tile_size(struct futhark_context_config *cfg, int size) {
  cfg->default_tile_size = size;
  cfg->default_tile_size_changed = 1;
}

void futhark_context_config_set_default_reg_tile_size(struct futhark_context_config *cfg, int size) {
  cfg->default_reg_tile_size = size;
}

void futhark_context_config_set_default_threshold(struct futhark_context_config *cfg, int size) {
  cfg->default_threshold = size;
}

int futhark_context_config_set_tuning_param(struct futhark_context_config *cfg,
                                            const char *param_name,
                                            size_t new_value) {
  for (int i = 0; i < cfg->num_tuning_params; i++) {
    if (strcmp(param_name, cfg->tuning_param_names[i]) == 0) {
      cfg->tuning_params[i] = new_value;
      return 0;
    }
  }
  if (strcmp(param_name, "default_group_size") == 0) {
    cfg->default_block_size = new_value;
    return 0;
  }
  if (strcmp(param_name, "default_num_groups") == 0) {
    cfg->default_grid_size = new_value;
    return 0;
  }
  if (strcmp(param_name, "default_threshold") == 0) {
    cfg->default_threshold = new_value;
    return 0;
  }
  if (strcmp(param_name, "default_tile_size") == 0) {
    cfg->default_tile_size = new_value;
    return 0;
  }
  if (strcmp(param_name, "default_reg_tile_size") == 0) {
    cfg->default_reg_tile_size = new_value;
    return 0;
  }
  return 1;
}

// A record of something that happened.
struct profiling_record {
  hipEvent_t *events; // Points to two events.
  const char *name;
};

struct futhark_context {
  struct futhark_context_config* cfg;
  int detail_memory;
  int debugging;
  int profiling;
  int profiling_paused;
  int logging;
  lock_t lock;
  char *error;
  lock_t error_lock;
  FILE *log;
  struct constants *constants;
  struct free_list free_list;
  int64_t peak_mem_usage_default;
  int64_t cur_mem_usage_default;
  // Uniform fields above.

  void* global_failure;
  void* global_failure_args;
  struct tuning_params tuning_params;
  // True if a potentially failing kernel has been enqueued.
  int32_t failure_is_an_option;
  int total_runs;
  long int total_runtime;
  int64_t peak_mem_usage_device;
  int64_t cur_mem_usage_device;
  struct program* program;

  hipDevice_t dev;
  int dev_id;
  hipModule_t module;
  hipStream_t stream;

  struct free_list gpu_free_list;

  size_t max_group_size;
  size_t max_grid_size;
  size_t max_tile_size;
  size_t max_threshold;
  size_t max_local_memory;
  size_t max_bespoke;

  size_t lockstep_width;

  struct profiling_record *profiling_records;
  int profiling_records_capacity;
  int profiling_records_used;

  struct builtin_kernels* kernels;
};

static int device_query(int dev_id, hipDeviceAttribute_t attr) {
  int val;
  HIP_SUCCEED_FATAL(hipDeviceGetAttribute(&val, attr, dev_id));
  return val;
}

static int function_query(hipFunction_t f, hipFunction_attribute attr) {
  int val;
  HIP_SUCCEED_FATAL(hipFuncGetAttribute(&val, attr, f));
  return val;
}

static int hip_device_setup(struct futhark_context *ctx) {
  struct futhark_context_config *cfg = ctx->cfg;
  int count, chosen = -1;
  hipDevice_t dev;

  HIP_SUCCEED_FATAL(hipGetDeviceCount(&count));
  if (count == 0) { return 1; }

  int num_device_matches = 0;

  for (int i = 0; i < count; i++) {
    hipDeviceProp_t prop;
    hipGetDeviceProperties(&prop, i);

    if (cfg->logging) {
      fprintf(ctx->log, "Device #%d: name=\"%s\"\n", i, prop.name);
    }

    if (strstr(prop.name, cfg->preferred_device) != NULL &&
        num_device_matches++ == cfg->preferred_device_num) {
      chosen = i;
      break;
    }
  }

  if (chosen == -1) { return 1; }

  if (cfg->logging) {
    fprintf(ctx->log, "Using device #%d\n", chosen);
  }

  ctx->dev_id = chosen;
  HIP_SUCCEED_FATAL(hipDeviceGet(&ctx->dev, ctx->dev_id));
  return 0;
}

static void hip_load_code_from_cache(struct futhark_context_config *cfg,
                                     const char *src,
                                     const char *opts[], size_t n_opts,
                                     struct cache_hash *h, const char *cache_fname,
                                     char **code, size_t *code_size) {
  if (cfg->logging) {
    fprintf(stderr, "Restoring cache from from %s...\n", cache_fname);
  }
  cache_hash_init(h);
  for (size_t i = 0; i < n_opts; i++) {
    cache_hash(h, opts[i], strlen(opts[i]));
  }
  cache_hash(h, src, strlen(src));
  errno = 0;
  if (cache_restore(cache_fname, h, (unsigned char**)code, code_size) != 0) {
    if (cfg->logging) {
      fprintf(stderr, "Failed to restore cache (errno: %s)\n", strerror(errno));
    }
  }
}

static void hip_size_setup(struct futhark_context *ctx) {
  struct futhark_context_config *cfg = ctx->cfg;
  if (cfg->default_block_size > ctx->max_group_size) {
    if (cfg->default_block_size_changed) {
      fprintf(stderr,
              "Note: Device limits default block size to %zu (down from %zu).\n",
              ctx->max_group_size, cfg->default_block_size);
    }
    cfg->default_block_size = ctx->max_group_size;
  }
  if (cfg->default_grid_size > ctx->max_grid_size) {
    if (cfg->default_grid_size_changed) {
      fprintf(stderr,
              "Note: Device limits default grid size to %zu (down from %zu).\n",
              ctx->max_grid_size, cfg->default_grid_size);
    }
    cfg->default_grid_size = ctx->max_grid_size;
  }
  if (cfg->default_tile_size > ctx->max_tile_size) {
    if (cfg->default_tile_size_changed) {
      fprintf(stderr,
              "Note: Device limits default tile size to %zu (down from %zu).\n",
              ctx->max_tile_size, cfg->default_tile_size);
    }
    cfg->default_tile_size = ctx->max_tile_size;
  }

  if (!cfg->default_grid_size_changed) {
    cfg->default_grid_size =
      (device_query(ctx->dev, hipDeviceAttributePhysicalMultiProcessorCount) *
       device_query(ctx->dev, hipDeviceAttributeMaxThreadsPerMultiProcessor))
      / cfg->default_block_size;
  }

  for (int i = 0; i < cfg->num_tuning_params; i++) {
    const char *size_class = cfg->tuning_param_classes[i];
    int64_t *size_value = &cfg->tuning_params[i];
    const char* size_name = cfg->tuning_param_names[i];
    int64_t max_value = 0, default_value = 0;

    if (strstr(size_class, "group_size") == size_class) {
      max_value = ctx->max_group_size;
      default_value = cfg->default_block_size;
    } else if (strstr(size_class, "num_groups") == size_class) {
      max_value = ctx->max_grid_size;
      default_value = cfg->default_grid_size;
      // XXX: as a quick and dirty hack, use twice as many threads for
      // histograms by default.  We really should just be smarter
      // about sizes somehow.
      if (strstr(size_name, ".seghist_") != NULL) {
        default_value *= 2;
      }
    } else if (strstr(size_class, "tile_size") == size_class) {
      max_value = ctx->max_tile_size;
      default_value = cfg->default_tile_size;
    } else if (strstr(size_class, "reg_tile_size") == size_class) {
      max_value = 0; // No limit.
      default_value = cfg->default_reg_tile_size;
    } else if (strstr(size_class, "threshold") == size_class) {
      // Threshold can be as large as it takes.
      default_value = cfg->default_threshold;
    } else {
      // Bespoke sizes have no limit or default.
    }

    if (*size_value == 0) {
      *size_value = default_value;
    } else if (max_value > 0 && *size_value > max_value) {
      fprintf(stderr, "Note: Device limits %s to %zu (down from %zu)\n",
              size_name, max_value, *size_value);
      *size_value = max_value;
    }
  }
}

static char* hiprtc_build(const char *src, const char *opts[], size_t n_opts,
                          char **code, size_t *code_size) {
  hiprtcProgram prog;
  char *problem = NULL;

  problem = HIPRTC_SUCCEED_NONFATAL(hiprtcCreateProgram(&prog, src, "futhark-hip", 0, NULL, NULL));

  if (problem) {
    return problem;
  }

  hiprtcResult res = hiprtcCompileProgram(prog, n_opts, opts);
  if (res != HIPRTC_SUCCESS) {
    size_t log_size;
    if (hiprtcGetProgramLogSize(prog, &log_size) == HIPRTC_SUCCESS) {
      char *log = (char*) malloc(log_size+1);
      log[log_size] = 0; // HIPRTC does not zero-terminate.
      if (hiprtcGetProgramLog(prog, log) == HIPRTC_SUCCESS) {
        problem = msgprintf("HIPRTC compilation failed.\n\n%s\n", log);
      } else {
        problem = msgprintf("Could not retrieve compilation log\n");
      }
      free(log);
    }
    return problem;
  }

  HIPRTC_SUCCEED_FATAL(hiprtcGetCodeSize(prog, code_size));
  *code = (char*) malloc(*code_size);
  HIPRTC_SUCCEED_FATAL(hiprtcGetCode(prog, *code));
  HIPRTC_SUCCEED_FATAL(hiprtcDestroyProgram(&prog));
  return NULL;
}

static void hiprtc_mk_build_options(struct futhark_context *ctx, const char *extra_opts[],
                                    char*** opts_out, size_t *n_opts) {
  int arch_set = 0, num_extra_opts;
  struct futhark_context_config *cfg = ctx->cfg;

  for (num_extra_opts = 0; extra_opts[num_extra_opts] != NULL; num_extra_opts++) {
    if (strstr(extra_opts[num_extra_opts], "--gpu-architecture")
        == extra_opts[num_extra_opts]) {
      arch_set = 1;
    }
  }

  size_t i = 0, n_opts_alloc = 20 + num_extra_opts + cfg->num_tuning_params;
  char **opts = (char**) malloc(n_opts_alloc * sizeof(char *));
  if (!arch_set) {
    hipDeviceProp_t props;
    HIP_SUCCEED_FATAL(hipGetDeviceProperties(&props, ctx->dev_id));
    opts[i++] = msgprintf("--gpu-architecture=%s", props.gcnArchName);
  }
  if (cfg->debugging) {
    opts[i++] = strdup("-G");
    opts[i++] = strdup("-lineinfo");
  }
  opts[i++] = msgprintf("-D%s=%d",
                        "max_group_size",
                        (int)ctx->max_group_size);
  for (int j = 0; j < cfg->num_tuning_params; j++) {
    opts[i++] = msgprintf("-D%s=%zu", cfg->tuning_param_vars[j],
                          cfg->tuning_params[j]);
  }
  opts[i++] = msgprintf("-DLOCKSTEP_WIDTH=%zu", ctx->lockstep_width);
  opts[i++] = msgprintf("-DMAX_THREADS_PER_BLOCK=%zu", ctx->max_group_size);

  for (int j = 0; extra_opts[j] != NULL; j++) {
    opts[i++] = strdup(extra_opts[j]);
  }

  opts[i++] = msgprintf("-DTR_BLOCK_DIM=%d", TR_BLOCK_DIM);
  opts[i++] = msgprintf("-DTR_TILE_DIM=%d", TR_TILE_DIM);
  opts[i++] = msgprintf("-DTR_ELEMS_PER_THREAD=%d", TR_ELEMS_PER_THREAD);

  *n_opts = i;
  *opts_out = opts;
}

static char* hip_module_setup(struct futhark_context *ctx,
                              const char *src,
                              const char *extra_opts[],
                              const char* cache_fname) {
  char *code = NULL;
  size_t code_size = 0;
  struct futhark_context_config *cfg = ctx->cfg;

  char **opts;
  size_t n_opts;
  hiprtc_mk_build_options(ctx, extra_opts, &opts, &n_opts);

  if (cfg->logging) {
    fprintf(stderr, "HIPRTC build options:\n");
    for (size_t j = 0; j < n_opts; j++) {
      fprintf(stderr, "\t%s\n", opts[j]);
    }
    fprintf(stderr, "\n");
  }

  struct cache_hash h;
  int loaded_code_from_cache = 0;
  if (cache_fname != NULL) {
    hip_load_code_from_cache(cfg, src, (const char**)opts, n_opts, &h, cache_fname, &code, &code_size);

    if (code != NULL) {
      if (cfg->logging) {
        fprintf(stderr, "Restored compiled code from cache; now loading module...\n");
      }
      if (hipModuleLoadData(&ctx->module, code) == hipSuccess) {
        if (cfg->logging) {
          fprintf(stderr, "Success!\n");
        }
        loaded_code_from_cache = 1;
      } else {
        if (cfg->logging) {
          fprintf(stderr, "Failed!\n");
        }
        free(code);
        code = NULL;
      }
    }
  }

  if (code == NULL) {
    char* problem = hiprtc_build(src, (const char**)opts, n_opts, &code, &code_size);
    if (problem != NULL) {
      return problem;
    }
  }

  if (!loaded_code_from_cache) {
    HIP_SUCCEED_FATAL(hipModuleLoadData(&ctx->module, code));
  }

  if (cache_fname != NULL && !loaded_code_from_cache) {
    if (cfg->logging) {
      fprintf(stderr, "Caching compiled code in %s...\n", cache_fname);
    }
    errno = 0;
    if (cache_store(cache_fname, &h, (const unsigned char*)code, code_size) != 0) {
      fprintf(stderr, "Failed to cache compiled code: %s\n", strerror(errno));
    }
  }

  for (size_t i = 0; i < n_opts; i++) {
    free((char *)opts[i]);
  }
  free(opts);
  free(code);

  return NULL;
}

static int tally_profiling_records(struct futhark_context *ctx,
                                   struct cost_centres* ccs) {
  hipError_t err;
  for (int i = 0; i < ctx->profiling_records_used; i++) {
    struct profiling_record record = ctx->profiling_records[i];

    float ms;
    if ((err = hipEventElapsedTime(&ms, record.events[0], record.events[1])) != hipSuccess) {
      return err;
    }

    if (ccs) {
      struct cost_centre c = {
        .name = record.name,
        .runs = 1,
        .runtime = ms*1000
      };
      cost_centres_add(ccs, c);
    }

    if ((err = hipEventDestroy(record.events[0])) != hipSuccess) {
      return 1;
    }
    if ((err = hipEventDestroy(record.events[1])) != hipSuccess) {
      return 1;
    }

    free(record.events);
  }

  ctx->profiling_records_used = 0;

  return 0;
}

static hipEvent_t* hip_get_events(struct futhark_context *ctx, const char* name) {
  if (ctx->profiling_records_used == ctx->profiling_records_capacity) {
    ctx->profiling_records_capacity *= 2;
    ctx->profiling_records =
      realloc(ctx->profiling_records,
              ctx->profiling_records_capacity *
              sizeof(struct profiling_record));
  }
  hipEvent_t *events = calloc(2, sizeof(hipEvent_t));
  hipEventCreate(&events[0]);
  hipEventCreate(&events[1]);
  ctx->profiling_records[ctx->profiling_records_used].events = events;
  ctx->profiling_records[ctx->profiling_records_used].name = name;
  ctx->profiling_records_used++;
  return events;
}

int futhark_context_sync(struct futhark_context* ctx) {
  HIP_SUCCEED_OR_RETURN(hipStreamSynchronize(ctx->stream));
  if (ctx->failure_is_an_option) {
    // Check for any delayed error.
    int32_t failure_idx;
    HIP_SUCCEED_OR_RETURN(hipMemcpyDtoH(&failure_idx,
                                        ctx->global_failure,
                                        sizeof(int32_t)));
    ctx->failure_is_an_option = 0;

    if (failure_idx >= 0) {
      // We have to clear global_failure so that the next entry point
      // is not considered a failure from the start.
      int32_t no_failure = -1;
      HIP_SUCCEED_OR_RETURN(hipMemcpyHtoD(ctx->global_failure,
                                          &no_failure,
                                          sizeof(int32_t)));

      int64_t args[max_failure_args+1];
      HIP_SUCCEED_OR_RETURN(hipMemcpyDtoH(&args,
                                          ctx->global_failure_args,
                                          sizeof(args)));

      ctx->error = get_failure_msg(failure_idx, args);

      return FUTHARK_PROGRAM_ERROR;
    }
  }
  return 0;
}

struct builtin_kernels* init_builtin_kernels(struct futhark_context* ctx);
void free_builtin_kernels(struct futhark_context* ctx, struct builtin_kernels* kernels);

int backend_context_setup(struct futhark_context* ctx) {
  ctx->profiling_records_capacity = 200;
  ctx->profiling_records_used = 0;
  ctx->profiling_records =
    malloc(ctx->profiling_records_capacity *
           sizeof(struct profiling_record));
  ctx->failure_is_an_option = 0;
  ctx->total_runs = 0;
  ctx->total_runtime = 0;
  ctx->peak_mem_usage_device = 0;
  ctx->cur_mem_usage_device = 0;

  HIP_SUCCEED_FATAL(hipInit(0));
  if (hip_device_setup(ctx) != 0) {
    futhark_panic(-1, "No suitable HIP device found.\n");
  }

  free_list_init(&ctx->gpu_free_list);

  ctx->max_local_memory = device_query(ctx->dev, hipDeviceAttributeMaxSharedMemoryPerBlock);
  ctx->max_group_size = device_query(ctx->dev, hipDeviceAttributeMaxThreadsPerBlock);
  ctx->max_grid_size = device_query(ctx->dev, hipDeviceAttributeMaxGridDimX);
  ctx->max_tile_size = sqrt(ctx->max_group_size);
  ctx->max_threshold = 0;
  ctx->max_bespoke = 0;
  // FIXME: in principle we should query hipDeviceAttributeWarpSize
  // from the device, which will provide 64 on AMD GPUs.
  // Unfortunately, we currently do nasty implicit intra-warp
  // synchronisation in codegen, which does not work when this is 64.
  // Once our codegen properly synchronises intra-warp operations, we
  // can use the actual hardware lockstep width instead.
  ctx->lockstep_width = 32;
  HIP_SUCCEED_FATAL(hipStreamCreate(&ctx->stream));
  hip_size_setup(ctx);
  ctx->error = hip_module_setup(ctx, ctx->cfg->program,
                                ctx->cfg->build_opts, ctx->cfg->cache_fname);

  if (ctx->error != NULL) {
    futhark_panic(1, "During HIP initialisation:\n%s\n", ctx->error);
  }

  int32_t no_error = -1;
  HIP_SUCCEED_FATAL(hipMalloc(&ctx->global_failure, sizeof(no_error)));
  HIP_SUCCEED_FATAL(hipMemcpyHtoD(ctx->global_failure, &no_error, sizeof(no_error)));
  // The +1 is to avoid zero-byte allocations.
  HIP_SUCCEED_FATAL(hipMalloc(&ctx->global_failure_args, sizeof(int64_t)*(max_failure_args+1)));

  if ((ctx->kernels = init_builtin_kernels(ctx)) == NULL) {
    return 1;
  }

  return 0;
}

void backend_context_teardown(struct futhark_context* ctx) {
  free_builtin_kernels(ctx, ctx->kernels);
  hipFree(ctx->global_failure);
  hipFree(ctx->global_failure_args);
  HIP_SUCCEED_FATAL(gpu_free_all(ctx));
  (void)tally_profiling_records(ctx, NULL);
  free(ctx->profiling_records);
  HIP_SUCCEED_FATAL(hipStreamDestroy(ctx->stream));
  HIP_SUCCEED_FATAL(hipModuleUnload(ctx->module));
}

// GPU ABSTRACTION LAYER

typedef hipFunction_t gpu_kernel;
typedef hipDeviceptr_t gpu_mem;

static void gpu_create_kernel(struct futhark_context *ctx,
                              gpu_kernel* kernel,
                              const char* name) {
  if (ctx->debugging) {
    fprintf(ctx->log, "Creating kernel %s.\n", name);
  }
  HIP_SUCCEED_FATAL(hipModuleGetFunction(kernel, ctx->module, name));
}

static void gpu_free_kernel(struct futhark_context *ctx,
                            gpu_kernel kernel) {
  (void)ctx;
  (void)kernel;
}

static int gpu_scalar_to_device(struct futhark_context* ctx,
                                gpu_mem dst, size_t offset, size_t size,
                                void *src) {
  hipEvent_t *pevents = NULL;
  if (ctx->profiling && !ctx->profiling_paused) {
    pevents = hip_get_events(ctx, "copy_scalar_to_dev");
    HIP_SUCCEED_FATAL(hipEventRecord(pevents[0], ctx->stream));
  }
  HIP_SUCCEED_OR_RETURN(hipMemcpyHtoD((unsigned char*)dst + offset, src, size));
  if (pevents != NULL) {
    HIP_SUCCEED_FATAL(hipEventRecord(pevents[1], ctx->stream));
  }
  return FUTHARK_SUCCESS;
}

static int gpu_scalar_from_device(struct futhark_context* ctx,
                                  void *dst,
                                  gpu_mem src, size_t offset, size_t size) {
  hipEvent_t *pevents = NULL;
  if (ctx->profiling && !ctx->profiling_paused) {
    pevents = hip_get_events(ctx, "copy_scalar_from_dev");
    HIP_SUCCEED_FATAL(hipEventRecord(pevents[0], ctx->stream));
  }
  HIP_SUCCEED_OR_RETURN(hipMemcpyDtoH(dst, (unsigned char*)src + offset, size));
  if (pevents != NULL) {
    HIP_SUCCEED_FATAL(hipEventRecord(pevents[1], ctx->stream));
  }
  return FUTHARK_SUCCESS;
}

static int gpu_memcpy(struct futhark_context* ctx,
                      gpu_mem dst, int64_t dst_offset,
                      gpu_mem src, int64_t src_offset,
                      int64_t nbytes) {
  hipEvent_t *pevents = NULL;
  if (ctx->profiling && !ctx->profiling_paused) {
    pevents = hip_get_events(ctx, "copy_dev_to_dev");
    HIP_SUCCEED_FATAL(hipEventRecord(pevents[0], ctx->stream));
  }
  HIP_SUCCEED_OR_RETURN(hipMemcpyWithStream((unsigned char*)dst+dst_offset, (unsigned char*)src+src_offset,
                                            nbytes, hipMemcpyDeviceToDevice ,ctx->stream));
  if (pevents != NULL) {
    HIP_SUCCEED_FATAL(hipEventRecord(pevents[1], ctx->stream));
  }
  return FUTHARK_SUCCESS;
}

static int memcpy_host2gpu(struct futhark_context* ctx, bool sync,
                           gpu_mem dst, int64_t dst_offset,
                           const unsigned char* src, int64_t src_offset,
                           int64_t nbytes) {
  if (nbytes > 0) {
    hipEvent_t* pevents = NULL;
    if (ctx->profiling && !ctx->profiling_paused) {
      pevents = hip_get_events(ctx, "copy_host_to_dev");
      HIP_SUCCEED_FATAL(hipEventRecord(pevents[0], ctx->stream));
    }
    if (sync) {
      HIP_SUCCEED_OR_RETURN
        (hipMemcpyHtoD((unsigned char*)dst + dst_offset,
                       (unsigned char*)src + src_offset, nbytes));
    } else {
      HIP_SUCCEED_OR_RETURN
        (hipMemcpyHtoDAsync((unsigned char*)dst + dst_offset,
                            (unsigned char*)src + src_offset,
                            nbytes, ctx->stream));
    }
    if (pevents != NULL) {
      HIP_SUCCEED_FATAL(hipEventRecord(pevents[1], ctx->stream));
    }
  }
  return FUTHARK_SUCCESS;
}

static int memcpy_gpu2host(struct futhark_context* ctx, bool sync,
                           unsigned char* dst, int64_t dst_offset,
                           gpu_mem src, int64_t src_offset,
                           int64_t nbytes) {
  if (nbytes > 0) {
    hipEvent_t* pevents = NULL;
    if (ctx->profiling && !ctx->profiling_paused) {
      pevents = hip_get_events(ctx, "copy_dev_to_host");
      HIP_SUCCEED_FATAL(hipEventRecord(pevents[0], ctx->stream));
    }
    if (sync) {
      HIP_SUCCEED_OR_RETURN
        (hipMemcpyDtoH(dst + dst_offset,
                       (unsigned char*)src + src_offset,
                       nbytes));
    } else {
      HIP_SUCCEED_OR_RETURN
        (hipMemcpyDtoHAsync(dst + dst_offset,
                            (unsigned char*)src + src_offset,
                            nbytes, ctx->stream));
    }
    if (sync &&
        ctx->failure_is_an_option &&
        futhark_context_sync(ctx) != 0) {
      return 1;
    }
  }
  return FUTHARK_SUCCESS;
}

static int gpu_launch_kernel(struct futhark_context* ctx,
                             gpu_kernel kernel, const char *name,
                             const int32_t grid[3],
                             const int32_t block[3],
                             unsigned int local_mem_bytes,
                             int num_args,
                             void* args[num_args],
                             size_t args_sizes[num_args]) {
  (void) args_sizes;
  int64_t time_start = 0, time_end = 0;
  if (ctx->logging) {
    fprintf(ctx->log,
            "Launching kernel %s with\n"
            "  grid=(%d,%d,%d)\n"
            "  block=(%d,%d,%d)\n"
            "  local memory=%d\n",
            name,
            grid[0], grid[1], grid[2],
            block[0], block[1], block[2],
            local_mem_bytes);
    time_start = get_wall_time();
  }

  hipEvent_t *pevents = NULL;
  if (ctx->profiling && !ctx->profiling_paused) {
    pevents = hip_get_events(ctx, name);
    HIP_SUCCEED_FATAL(hipEventRecord(pevents[0], ctx->stream));
  }

  HIP_SUCCEED_OR_RETURN
    (hipModuleLaunchKernel(kernel,
                           grid[0], grid[1], grid[2],
                           block[0], block[1], block[2],
                           local_mem_bytes, ctx->stream,
                           args, NULL));

  if (pevents != NULL) {
    HIP_SUCCEED_FATAL(hipEventRecord(pevents[1], ctx->stream));
  }

  if (ctx->debugging) {
    HIP_SUCCEED_FATAL(hipStreamSynchronize(ctx->stream));
    time_end = get_wall_time();
    long int time_diff = time_end - time_start;
    fprintf(ctx->log, "  runtime: %ldus\n\n", time_diff);
  }

  return FUTHARK_SUCCESS;
}

static int gpu_alloc_actual(struct futhark_context *ctx, size_t size, gpu_mem *mem_out) {
  hipError_t res = hipMalloc(mem_out, size);
  if (res == hipErrorOutOfMemory) {
    return FUTHARK_OUT_OF_MEMORY;
  }
  HIP_SUCCEED_OR_RETURN(res);
  return FUTHARK_SUCCESS;
}

static int gpu_free_actual(struct futhark_context *ctx, gpu_mem mem) {
  (void)ctx;
  HIP_SUCCEED_OR_RETURN(hipFree(mem));
  return FUTHARK_SUCCESS;
}

// End of backends/hip.h.