Example defining an intrinsic in the Numba CUDA target

cuda_intrinsic_example.py

from numba.cuda.extending import intrinsic
from llvmlite import ir

@intrinsic
def cuda_clock64(typingctx):
    sig = types.uint64()

    def codegen(context, builder, sig, args):
        function_type = ir.FunctionType(ir.IntType(64), []) 
        instruction = "mov.u64 $0, %clock64;"
        clock64 = ir.InlineAsm(function_type, instruction, "=l",
                               side_effect=True)
        return builder.call(clock64, []) 

    return sig, codegen