Graham Markall gmarkall
gmarkall
/ cupy_pycuda_arrays.py
Created
February 19, 2021 11:59
Calling PyCUDA kernels on CuPy arrays using the CUDA array interface
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| 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 |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| 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
/ mcjit.cpp
Created
February 10, 2021 13:03
— forked from gligneul/mcjit.cpp
LLVM MCJIT Code Samples (Working!)
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| /* | |
| * 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
/ fix-numba-051-2.patch
Created
August 18, 2020 19:59
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| 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
/ fix-numba-051.patch
Created
August 18, 2020 19:58
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| 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
/ compile_to_ptx.py
Created
August 7, 2020 11:23
Numba runtime/driver calls when compiling to PTX
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| 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] | |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| $ 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
/ 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
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| # 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 |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| 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
/ .bash_profile
Created
December 7, 2015 15:34
ssh-agent setup on windows with git bash 2.6.3
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| # 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 |