diff --git a/numba_cuda/numba/cuda/debuginfo.py b/numba_cuda/numba/cuda/debuginfo.py
index 81f91f5ff..65136119d 100644
--- a/numba_cuda/numba/cuda/debuginfo.py
+++ b/numba_cuda/numba/cuda/debuginfo.py
@@ -59,6 +59,33 @@ def _var_type(self, lltype, size, datamodel=None):
         # For other cases, use upstream Numba implementation
         return super()._var_type(lltype, size, datamodel=datamodel)
 
+    def _di_subroutine_type(self, line, function, argmap):
+        # The function call conv needs encoding.
+        llfunc = function
+        md = []
+
+        # Create metadata type for return value
+        if len(llfunc.args) > 0:
+            lltype = llfunc.args[0].type
+            size = self.cgctx.get_abi_sizeof(lltype)
+            mdtype = self._var_type(lltype, size, datamodel=None)
+            md.append(mdtype)
+
+        # Create metadata type for arguments
+        for idx, (name, nbtype) in enumerate(argmap.items()):
+            datamodel = self.cgctx.data_model_manager[nbtype]
+            lltype = self.cgctx.get_value_type(nbtype)
+            size = self.cgctx.get_abi_sizeof(lltype)
+            mdtype = self._var_type(lltype, size, datamodel=datamodel)
+            md.append(mdtype)
+
+        return self.module.add_debug_info(
+            "DISubroutineType",
+            {
+                "types": self.module.add_metadata(md),
+            },
+        )
+
     def mark_variable(
         self,
         builder,
diff --git a/numba_cuda/numba/cuda/target.py b/numba_cuda/numba/cuda/target.py
index 5a49bffb4..b97bf1203 100644
--- a/numba_cuda/numba/cuda/target.py
+++ b/numba_cuda/numba/cuda/target.py
@@ -290,7 +290,16 @@ def get_ufunc_info(self, ufunc_key):
 
 
 class CUDACallConv(MinimalCallConv):
-    pass
+    def decorate_function(self, fn, args, fe_argtypes, noalias=False):
+        """
+        Set names and attributes of function arguments.
+        """
+        assert not noalias
+        arginfo = self._get_arg_packer(fe_argtypes)
+        # Do not prefix "arg." on argument name, so that nvvm compiler
+        # can track debug info of argument more accurately
+        arginfo.assign_names(self.get_arguments(fn), args)
+        fn.args[0].name = ".ret"
 
 
 class CUDACABICallConv(BaseCallConv):
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py
index 6b0065201..399c09d5a 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py
@@ -310,6 +310,23 @@ def test_kernel_args_types_dump(self):
             with captured_stdout():
                 self._test_kernel_args_types()
 
+    def test_kernel_args_names(self):
+        sig = (types.int32,)
+
+        @cuda.jit("void(int32)", debug=True, opt=False)
+        def f(x):
+            z = x  # noqa: F841
+
+        llvm_ir = f.inspect_llvm(sig)
+
+        # Verify argument name is not prefixed with "arg."
+        pat = r"define void @.*\(i32 %\"x\"\)"
+        match = re.compile(pat).search(llvm_ir)
+        self.assertIsNotNone(match, msg=llvm_ir)
+        pat = r"define void @.*\(i32 %\"arg\.x\"\)"
+        match = re.compile(pat).search(llvm_ir)
+        self.assertIsNone(match, msg=llvm_ir)
+
     def test_llvm_dbg_value(self):
         sig = (types.int32, types.int32)