Skip to content

Instantly share code, notes, and snippets.

View gmarkall's full-sized avatar

Graham Markall gmarkall

239 followers · 16 following
View GitHub Profile
@gmarkall
gmarkall / cupy_pycuda_arrays.py
Created February 19, 2021 11:59
Calling PyCUDA kernels on CuPy arrays using the CUDA array interface
import pycuda.driver as cuda
import pycuda.autoinit # noqa
from collections import namedtuple
from pycuda.compiler import SourceModule
import cupy as cp
# PyCUDA will try to get a pointer to data from an object it doesn't recognise
@gmarkall
gmarkall / cuda_pipeline.diff
Created February 18, 2021 15:42
Adding a CUDA pipeline to Numba
diff --git a/numba/core/compiler.py b/numba/core/compiler.py
index a0c03fdb3..e47270a2f 100644
--- a/numba/core/compiler.py
+++ b/numba/core/compiler.py
@@ -28,7 +28,8 @@ from numba.core.typed_passes import (NopythonTypeInference, AnnotateTypes,
NopythonRewrites, PreParforPass,
ParforPass, DumpParforDiagnostics,
IRLegalization, NoPythonBackend,
- InlineOverloads, PreLowerStripPhis)
+ InlineOverloads, PreLowerStripPhis,
@gmarkall
gmarkall / mcjit.cpp
Created February 10, 2021 13:03 — forked from gligneul/mcjit.cpp
LLVM MCJIT Code Samples (Working!)
/*
* To compile, execute on terminal:
* g++ -o mcjit mcjit.cpp `llvm-config --cxxflags --ldflags --libs all --system-libs`
*/
#include <iostream>
#include <memory>
#include <llvm/ADT/StringRef.h>
#include <llvm/ExecutionEngine/ExecutionEngine.h>
@gmarkall
gmarkall / fix-numba-051-2.patch
Created August 18, 2020 19:59
diff --git a/numba/cuda/cudadrv/nvvm.py b/numba/cuda/cudadrv/nvvm.py
index 24569f99b..505e6797f 100644
--- a/numba/cuda/cudadrv/nvvm.py
+++ b/numba/cuda/cudadrv/nvvm.py
@@ -272,29 +272,39 @@ data_layout = {
default_data_layout = data_layout[tuple.__itemsize__ * 8]
+_supported_cc = None
@gmarkall
gmarkall / fix-numba-051.patch
Created August 18, 2020 19:58
diff --git a/numba/cuda/cudadrv/nvvm.py b/numba/cuda/cudadrv/nvvm.py
index 24569f99b..505e6797f 100644
--- a/numba/cuda/cudadrv/nvvm.py
+++ b/numba/cuda/cudadrv/nvvm.py
@@ -272,29 +272,39 @@ data_layout = {
default_data_layout = data_layout[tuple.__itemsize__ * 8]
+_supported_cc = None
@gmarkall
gmarkall / compile_to_ptx.py
Created August 7, 2020 11:23
Numba runtime/driver calls when compiling to PTX
from numba import cuda, types
import time
def add(r, x, y):
i = cuda.grid(1)
if i < len(r):
r[i] = x[i] + y[i]
@gmarkall
gmarkall / output.txt
Created August 7, 2020 11:16
Numba CUDA kernel compile and launch traces
$ NUMBA_CUDA_LOG_LEVEL=DEBUG python repro.py
== CUDA [182] DEBUG -- call runtime api: cudaRuntimeGetVersion
Define a kernel
Copy to device
== CUDA [200] INFO -- init
== CUDA [200] DEBUG -- call driver api: cuInit
== CUDA [200] DEBUG -- call driver api: cuCtxGetCurrent
== CUDA [200] DEBUG -- call driver api: cuCtxGetCurrent
== CUDA [200] DEBUG -- call driver api: cuDeviceGetCount
== CUDA [200] DEBUG -- call driver api: cuDeviceGet
@gmarkall
gmarkall / extending_cuda.py
Created June 16, 2020 15:15
Extending the CUDA target to support an Interval type - based on https://numba.pydata.org/numba-doc/latest/extending/interval-example.html with some modifications
# From Interval example docs - typing
from numba import types
from numba.core.extending import typeof_impl, type_callable
# Data model
from numba.core.extending import models, register_model
# Lowering
@gmarkall
gmarkall / dtest.py
Created January 4, 2019 11:06
import hashlib
d = {}
for i in range(1200000):
k = hashlib.md5('%x' % i).digest()
d[k] = i
#for k in d.keys():
# print k
@gmarkall
gmarkall / .bash_profile
Created December 7, 2015 15:34
ssh-agent setup on windows with git bash 2.6.3
# Based on http://mah.everybody.org/docs/ssh
# set environment variables if user's agent already exists
[ -z "$SSH_AUTH_SOCK" ] && SSH_AUTH_SOCK=$(ls -l /tmp/ssh-*/agent.* 2> /dev/null | grep $(whoami) | awk '{print $9}')
[ -z "$SSH_AGENT_PID" -a "$SSH_AUTH_SOCK" ] && [ -n `echo $SSH_AUTH_SOCK | cut -d. -f2` ] && SSH_AGENT_PID=$((`echo $SSH_AUTH_SOCK | cut -d. -f2` + 1))
[ -n "$SSH_AUTH_SOCK" ] && export SSH_AUTH_SOCK
[ -n "$SSH_AGENT_PID" ] && export SSH_AGENT_PID
# start agent if necessary
if [ -z $SSH_AGENT_PID ] && [ -z $SSH_TTY ]; then # if no agent & not in ssh