Skip to content

Instantly share code, notes, and snippets.

View jaredhoberock's full-sized avatar

Jared Hoberock jaredhoberock

318 followers · 10 following
  • NVIDIA
  • Kansas
View GitHub Profile
@jaredhoberock
jaredhoberock / location_of_delete_keyword.cpp
Created August 18, 2023 04:58
Lexing for delete
SourceLocation location_of_delete_keyword(FunctionDecl* FD, SourceManager& SM, const LangOptions& LO)
{
// getSourceRange() of a CXXMethodDecl seems to include the "= delete"
if(isa<CXXMethodDecl>(FD)) return FD->getSourceRange().getEnd();
// we need to lex to find the end of "delete" for functions
// that are not methods because their SourceRange omits
// the trailing "= delete"
SourceLocation loc = FD->getBeginLoc();
SourceLocation result = FD->getEndLoc();
@jaredhoberock
jaredhoberock / neo_executor.cpp
Created August 2, 2018 16:07
Experimental polymorphic executor which exposes a single .execute() function
// $ clang-5.0 -std=c++17 -Iexecutors-impl/include -lstdc++ neo_executor.cpp
#include <experimental/execution>
#include <type_traits>
#include <iostream>
template<class T>
struct __is_bulk_task
{
template<class U = T,
class = decltype(std::declval<U>().function()),
@jaredhoberock
jaredhoberock / reduce.hpp
Created May 29, 2018 15:27
Possible implementation of std::reduce via executors
// Copyright (c) 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
// documentation and/or other materials provided with the distribution.
@jaredhoberock
jaredhoberock / eager_executor.hpp
Last active October 13, 2016 18:29
Example of an executor whose asynchronous operations always return a ready future
template<class T>
class always_ready_future
{
public:
always_ready_future(T&& value)
: value_(std::move(value))
{}
// never blocks
void wait(){}
@jaredhoberock
jaredhoberock / repro.cpp
Created February 28, 2015 02:32
Empty tuples
#include <type_traits>
#include <cassert>
template<int i, class Whence, class T, bool = std::is_empty<T>::value>
struct tuple_leaf
{
tuple_leaf(T x) : val(x) {}
T& get() { return val; }
@jaredhoberock
jaredhoberock / sort_host_device.cu
Created July 15, 2014 02:39
CUB __host__ __device__ sort
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__>= 350 && defined(__CUDACC_RDC__))
# define CUB_CDP 1
#endif
#if defined(__CUDACC__)
# if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__>= 350 && defined(__CUDACC_RDC__))
# define HAS_CUDART 1
# else
# define HAS_CUDART 0
# endif
@jaredhoberock
jaredhoberock / Particles.cu
Created June 7, 2012 20:59 — forked from dangets/Particles.cu
CUDA Thrust Structure of Arrays reference as normal structure (working)
template<typename Vector>
struct Particles {
typedef typename Vector::value_type T;
typedef thrust::zip_iterator<
thrust::tuple<
typename Vector::iterator,
typename Vector::iterator,
@jaredhoberock
jaredhoberock / gist:2561193
Created April 30, 2012 18:51
Thrust article listings
// Figure 16.1
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/generate.h>
#include <thrust/sort.h>
#include <thrust/copy.h>
#include <cstdlib>
int main()
@jaredhoberock
jaredhoberock / merge_by_key_example.cu
Created April 6, 2012 06:11
How to build merge_by_key using thrust::merge and thrust::zip_iterator
#include <thrust/pair.h>
#include <thrust/tuple.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/functional.h>
#include <thrust/merge.h>
#include <iostream>
#include <algorithm>
// compares the 0th elements of two tuples
template<typename Compare>