[NVVM IR] NVVM IR frontend support for cuda core

[abhilash1910]

Abstract : The cuda/bindings backend of Cuda python has NVVM support through libnvvm api . However the frontend of cuda python does not support nvvm ir as input source. Since cuda python allows users to leverage a "pythonic dsl" format for writing the host code (taking care of launch parameters etc), it makes sense to also allow NVVM IR as an alternative input to the already included list of inputs {ptx, c++, lto ir} etc.

Technicalities : The integration seems straightforward as hooking the cuda core frontend with the already existing nvvm ir backend inside cuda bindings. However for the plugin for NVVM, the current proposal dumps only the NVVM IR compiled Program - i.e there is no symbolic linking or processing for downstream tasks. As NVVM currently flows through the PTX pathway, the current design does not add that module as it requires additional design decisions.

Current proposition flow :
Assuming a NVVM IR ASCII string as input as :

nvvm_program= """target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-i128:128:128-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
target triple = "nvptx64-nvidia-cuda"

define void @simple_kernel(i32* %output) {
entry:
  %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
  %bid = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()  
  %bdim = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
  %idx = add i32 %tid, mul i32 %bid, %bdim
  
  %idx64 = zext i32 %idx to i64
  %ptr = getelementptr inbounds i32, i32* %output, i64 %idx64
  store i32 %idx, i32* %ptr, align 4
  ret void
}

declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() nounwind readnone
declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() nounwind readnone
declare i32 @llvm.nvvm.read.ptx.sreg.ntid.x() nounwind readnone

!nvvm.annotations = !{!0}
!0 = !{void (i32*)* @simple_kernel, !"kernel", i32 1}

!nvvmir.version = !{!1}
!1 = !{i32 2, i32 0, i32 3, i32 0}
"""
the invocation to the framework would be  like so:

``` python
options = ProgramOptions(
        name="simple_nvvm_program",
        arch="sm_70",              # Target specific architecture
        debug=False,               # No debug info for performance
        device_code_optimize=True  # Enable optimizations
    )
program = Program(nvvm_program, "nvvm", options=options)
# This program can be then be compiled to ptx
object_code = program.compile("ptx")

Extractions : Currently the proposal does not add anything which relates to the mapping of symbols and follows the existing PTX pathway. NVRTC pathway has a symbolic mapping like so:

if name_expressions:
    for n in name_expressions:
        nvrtc.nvrtcAddNameExpression(self._mnff.handle, n.encode())
# ... compile ...
symbol_mapping = {}
if name_expressions:
    for n in name_expressions:
        symbol_mapping[n] = nvrtc.nvrtcGetLoweredName(self._mnff.handle, n.encode())

return ObjectCode._init(data, target_type, 
                       symbol_mapping=symbol_mapping, ...)  # ← mapping

Even in the case of PTX or LTO IR backend, users have to manually specify the symbolic linkages:

 mod_obj = ObjectCode.from_ptx(ptx, symbol_mapping=sym_map)

The current proposition for NVVM IR is same as that of the PTX counterpart, as eventually the NVVM IR boils down to PTX .
However it would be nice to have design which would bypass the ptx counterpart for symbolic execution.

There is a draft local PR which will be linked here for the addition.
Edit: #907
cc @leofang

[leofang]

Thanks a lot, Abhilash! This is tracked in #452.