brycelelbach

  Before long they had reached the edge of the settlement and made their way
beside a sluggish stream through bare-twigged scrubby bushes. Occasionally
they would hear a harsh croak or a splash as some amphibian was disturbed, but
the only creature they saw was a toad as big as Will’s foot, which could only
flop in a pain-filled sideways heave as if it were horribly injured. It lay
across the path, trying to move out of the way and looking at them as if it
knew they meant to hurt it.
  “It would be merciful to kill it,” said Tialys.
  “How do you know?” said Lyra. “It might still like being alive, in spite of
everything.”

brycelelbach

#include <cassert>

int current_device()
{
  int device = 0;
  cudaError_t const error = cudaGetDevice(&device);
  assert(cudaSuccess == error);
  return device;
}

brycelelbach

Blind the Submissions

Motivation

The first round of reviews of submissions to technical conferences should be double-blind (e.g. reviewers don't know who the submitter is).

Non-double-blind submissions:

• Contribute to Hero Culture: Hero culture is the tendency within technical

brycelelbach

/******************************************************************************
 * Copyright (c) 2011, Duane Merrill.  All rights reserved.
 * Copyright (c) 2011-2018, NVIDIA CORPORATION.  All rights reserved.
 *
 * Redistribution and use in source and binary forms, with or without
 * modification, are permitted provided that the following conditions are met:
 *     * Redistributions of source code must retain the above copyright
 *       notice, this list of conditions and the following disclaimer.
 *     * Redistributions in binary form must reproduce the above copyright
 *       notice, this list of conditions and the following disclaimer in the

brycelelbach

// This is how we run libc++ tests on the GPU without modification.
// We force include this header into each test with `-include`.

__host__ __device__
int fake_main(int, char**);

__global__
void fake_main_kernel(int * ret)
{
    *ret = fake_main(0, NULL);

brycelelbach

// I have this code:

struct thread_group {
private:
  std::vector<std::thread> members;

public:
  thread_group(thread_group const&) = delete;
  thread_group& operator=(thread_group const&) = delete;

brycelelbach

// Sort the sequence of integers by the Nth bit.
template <typename ExecutionPolicy,
          std::ranges::random_access_range InputRange, random_access_iterator OutputIt>
  requires std::integral<typename std::ranges_value_t<InputRange>>
unique_future<std::uint64_t> async_radix_sort_pass(ExecutionPolicy&& exec,
                                                   InputRange input, OutputRange output,
                                                   std::uint64_t bit)
{
  auto const elements = std::distance(input);

brycelelbach

template <typename InputIt, typename OutputIt>
OutputIt
radix_sort_split(InputIt first, InputIt last, OutputIt output, std::uint64_t bit)
{
  std::vector<std::uint64_t> e(std::distance(first, last));

  // Count 0s.
  std::transform(first, last, e.begin(),
                 [=] (auto t) { return !(t & (1 << bit)); });

brycelelbach

template <typename InputIterator, typename OutputIterator, typename T, typename BinaryOp>
OutputIterator exclusive_scan(InputIterator first, InputIterator last,
                              OutputIterator result, T init, BinaryOp op)
{
  if (first != last) {
    T saved = init;
    do {
      init = op(init, *first);
      *result = saved;
      saved = init;

brycelelbach

template <typename InputIt, typename OutputIt, typename BinaryOp, typename T, typename Size>
unique_future<OutputIt>
async_inclusive_scan(InputIt first, InputIt last, OutputIt output,BinaryOp op, T init, Size chunk_size)
{
  Size const elements = std::distance(first, last);
  Size const chunks   = (1 + ((elements - 1) / chunk_size)); // Round up.

  std::vector<unique_future<T>> sweep;
  sweep.reserve(chunks);