Skip to content

Instantly share code, notes, and snippets.

@gmarkall
Created August 7, 2020 11:23
Show Gist options
  • Save gmarkall/e2071a844e8db256dc9f53f11f3a1db5 to your computer and use it in GitHub Desktop.
Save gmarkall/e2071a844e8db256dc9f53f11f3a1db5 to your computer and use it in GitHub Desktop.
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]
def sub(r, x, y):
i = cuda.grid(1)
if i < len(r):
r[i] = x[i] - y[i]
arg_types = (types.float64[::1], types.float64[::1], types.float64[::1])
cc = (7, 5)
print("Compile to PTX")
cuda.compile_ptx(add, arg_types)
print("Compile to PTX for current device")
cuda.compile_ptx_for_current_device(sub, arg_types)
$ NUMBA_CUDA_LOG_LEVEL=DEBUG python compile_to_ptx.py
== CUDA [181] DEBUG -- call runtime api: cudaRuntimeGetVersion
Compile to PTX
Compile to PTX for current device
== CUDA [292] INFO -- init
== CUDA [292] DEBUG -- call driver api: cuInit
== CUDA [292] DEBUG -- call driver api: cuCtxGetCurrent
== CUDA [292] DEBUG -- call driver api: cuDeviceGetCount
== CUDA [292] DEBUG -- call driver api: cuDeviceGet
== CUDA [292] DEBUG -- call driver api: cuDeviceComputeCapability
== CUDA [292] DEBUG -- call driver api: cuDeviceGetName
== CUDA [293] DEBUG -- call driver api: cuDevicePrimaryCtxRetain
== CUDA [398] DEBUG -- call driver api: cuCtxPushCurrent_v2
== CUDA [398] DEBUG -- call driver api: cuMemGetInfo_v2
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment