[FEA] Support for simpler ways to inline PTX

[gmarkall]

Presently PTX can be inlined, but it requires writing NVVM IR with llvmlite inside a Numba extension (either with typing and a lowering function that generates an ir.InlineAsm instruction, or an overload and an intrinsic with the necessary code generation).

It would be nicer to be able to write inline PTX more simply. There are two possible ways this could be done:

• Through Pythonic intrinsics for PTX instructions, if these were to exist.
• Through a simpler API that can be used directly in a kernel.

For the latter option, the usage could look like:

@cuda.jit
def f(r, x):
    arg = x[0]
    result = inline_ptx("tanh.approx.f32 $0, $1;", "=f,f", (arg,))
    r[0] = result

This mimics the CUDA C++ API for inline PTX, where the assembly snippet, constraints, and arguments all need to be provided.

cc @leofang @oleksandr-pavlyk @benhg

[leofang]

For the former option:

• Through Pythonic intrinsics for PTX instructions, if these were to exist.

This is one such example, although the interfaces would require a careful design to cover all possible PTX instructions, operands and directives.

@cuda.jit
def f(r, x):
    arg = x[0]
    # maps to tanh.approx.f32 if arg.dtype is float32 and likewise for all supported dtypes,
    # otherwise raise a compile-time error
    result = tanh.approx(arg) 
    r[0] = result

[kkraus14]

• Through Pythonic intrinsics for PTX instructions, if these were to exist.
• Through a simpler API that can be used directly in a kernel.

I don't think these are mutually exclusive and we should support both.

For better Pythonic intrinsics for PTX instructions, I hope we can leverage NumbAST + libcu++ PTX API (https://nvidia.github.io/cccl/libcudacxx/ptx_api.html) to do most of the heavy lifting for us.

There will always be reasons people want to use inline PTX directly and we should support it the same way it's available in CUDA C++.

[leofang]

xref: NVIDIA/cccl#6459