diff --git a/numba_cuda/numba/cuda/cgutils.py b/numba_cuda/numba/cuda/cgutils.py
index b07a83031..fee2489bc 100644
--- a/numba_cuda/numba/cuda/cgutils.py
+++ b/numba_cuda/numba/cuda/cgutils.py
@@ -92,7 +92,7 @@ def copy_struct(dst, src, repl=None):
     return dst
 
 
-class _StructProxy(object):
+class _StructProxy:
     """
     Creates a `Structure` like interface that is constructed with information
     from DataModel instance.  FE type must have a data model that is a
@@ -171,7 +171,7 @@ def __setattr__(self, field, value):
         Store the LLVM *value* into the named *field*.
         """
         if field.startswith("_"):
-            return super(_StructProxy, self).__setattr__(field, value)
+            return super().__setattr__(field, value)
         self[self._datamodel.get_field_position(field)] = value
 
     def __getitem__(self, index):
@@ -269,7 +269,7 @@ def _cast_member_from_value(self, index, val):
         return model.as_data(self._builder, val)
 
 
-class Structure(object):
+class Structure:
     """
     A high-level object wrapping a alloca'ed LLVM structure, including
     named fields and attribute access.
@@ -329,7 +329,7 @@ def __setattr__(self, field, value):
         Store the LLVM *value* into the named *field*.
         """
         if field.startswith("_"):
-            return super(Structure, self).__setattr__(field, value)
+            return super().__setattr__(field, value)
         self[self._namemap[field]] = value
 
     def __getitem__(self, index):
diff --git a/numba_cuda/numba/cuda/compiler.py b/numba_cuda/numba/cuda/compiler.py
index 5b031d06b..511d28eb4 100644
--- a/numba_cuda/numba/cuda/compiler.py
+++ b/numba_cuda/numba/cuda/compiler.py
@@ -117,7 +117,7 @@ def run_frontend(func, inline_closures=False, emit_dels=False):
     return func_ir
 
 
-class DefaultPassBuilder(object):
+class DefaultPassBuilder:
     """
     This is the default pass builder, it contains the "classic" default
     pipelines as pre-canned PassManager instances:
diff --git a/numba_cuda/numba/cuda/core/analysis.py b/numba_cuda/numba/cuda/core/analysis.py
index b849891a8..f3a86b1ba 100644
--- a/numba_cuda/numba/cuda/core/analysis.py
+++ b/numba_cuda/numba/cuda/core/analysis.py
@@ -555,7 +555,7 @@ def prune_by_predicate(branch, pred, blk):
         taken = do_prune(take_truebr, blk)
         return True, taken
 
-    class Unknown(object):
+    class Unknown:
         pass
 
     def resolve_input_arg_const(input_arg_idx):
diff --git a/numba_cuda/numba/cuda/core/annotations/type_annotations.py b/numba_cuda/numba/cuda/core/annotations/type_annotations.py
index 8c426322c..cb18f7db3 100644
--- a/numba_cuda/numba/cuda/core/annotations/type_annotations.py
+++ b/numba_cuda/numba/cuda/core/annotations/type_annotations.py
@@ -42,7 +42,7 @@ def avail(self):
         return bool(self.lines)
 
 
-class TypeAnnotation(object):
+class TypeAnnotation:
     # func_data dict stores annotation data for all functions that are
     # compiled. We store the data in the TypeAnnotation class since a new
     # TypeAnnotation instance is created for each function that is compiled.
diff --git a/numba_cuda/numba/cuda/core/base.py b/numba_cuda/numba/cuda/core/base.py
index 617ced951..f09e512f9 100644
--- a/numba_cuda/numba/cuda/core/base.py
+++ b/numba_cuda/numba/cuda/core/base.py
@@ -30,7 +30,7 @@
 void_ptr = GENERIC_POINTER
 
 
-class OverloadSelector(object):
+class OverloadSelector:
     """
     An object matching an actual signature against a registry of formal
     signatures and choosing the best candidate, if any.
@@ -151,7 +151,7 @@ def append(self, value, sig):
         self._cache.clear()
 
 
-class BaseContext(object):
+class BaseContext:
     """
 
     Notes on Structure
@@ -1271,7 +1271,7 @@ def get_ufunc_info(self, ufunc_key):
         raise NotImplementedError(f"{self} does not support ufunc")
 
 
-class _wrap_impl(object):
+class _wrap_impl:
     """
     A wrapper object to call an implementation function with some predefined
     (context, signature) arguments.
@@ -1302,7 +1302,7 @@ def _has_loc(fn):
     return "loc" in sig.parameters
 
 
-class _wrap_missing_loc(object):
+class _wrap_missing_loc:
     def __init__(self, fn):
         self.func = fn  # store this to help with debug
 
diff --git a/numba_cuda/numba/cuda/core/boxing.py b/numba_cuda/numba/cuda/core/boxing.py
index 3269ba73b..eb3b0c6ee 100644
--- a/numba_cuda/numba/cuda/core/boxing.py
+++ b/numba_cuda/numba/cuda/core/boxing.py
@@ -673,7 +673,7 @@ def box_list(typ, val, c):
     return c.builder.load(res)
 
 
-class _NumbaTypeHelper(object):
+class _NumbaTypeHelper:
     """A helper for acquiring `numba.typeof` for type checking.
 
     Usage
diff --git a/numba_cuda/numba/cuda/core/bytecode.py b/numba_cuda/numba/cuda/core/bytecode.py
index 13ad54e84..32dc4e353 100644
--- a/numba_cuda/numba/cuda/core/bytecode.py
+++ b/numba_cuda/numba/cuda/core/bytecode.py
@@ -70,7 +70,7 @@ def _as_opcodes(seq):
 HAVE_ARGUMENT = dis.HAVE_ARGUMENT
 
 
-class ByteCodeInst(object):
+class ByteCodeInst:
     """
     Attributes
     ----------
@@ -271,7 +271,7 @@ def _patched_opargs(bc_stream):
         yield offset + _FIXED_OFFSET, opcode, arg, nextoffset + _FIXED_OFFSET
 
 
-class ByteCodeIter(object):
+class ByteCodeIter:
     def __init__(self, code):
         self.code = code
         self.iter = iter(_patched_opargs(_unpack_opargs(self.code.co_code)))
@@ -298,7 +298,7 @@ def read_arg(self, size):
         return buf
 
 
-class _ByteCode(object):
+class _ByteCode:
     """
     The decoded bytecode of a function, and related information.
     """
diff --git a/numba_cuda/numba/cuda/core/byteflow.py b/numba_cuda/numba/cuda/core/byteflow.py
index 8f8959caf..46eef06f9 100644
--- a/numba_cuda/numba/cuda/core/byteflow.py
+++ b/numba_cuda/numba/cuda/core/byteflow.py
@@ -51,7 +51,7 @@ class CALL_INTRINSIC_1_Operand(Enum):
 
 
 @total_ordering
-class BlockKind(object):
+class BlockKind:
     """Kinds of block to make related code safer than just `str`."""
 
     _members = frozenset(
@@ -88,7 +88,7 @@ def __repr__(self):
         return "BlockKind({})".format(self._value)
 
 
-class Flow(object):
+class Flow:
     """Data+Control Flow analysis.
 
     Simulate execution to recover dataflow and controlflow information.
@@ -353,7 +353,7 @@ def _is_null_temp_reg(reg):
     return reg.startswith("$null$")
 
 
-class TraceRunner(object):
+class TraceRunner:
     """Trace runner contains the states for the trace and the opcode dispatch."""
 
     def __init__(self, debug_filename):
@@ -2118,7 +2118,7 @@ def op_LOAD_COMMON_CONSTANT(self, state, inst):
 
 
 @total_ordering
-class _State(object):
+class _State:
     """State of the trace"""
 
     def __init__(self, bytecode, pc, nstack, blockstack, nullvals=()):
@@ -2487,7 +2487,7 @@ def get_function_attributes(self, make_func_res):
 Edge = namedtuple("Edge", ["pc", "stack", "blockstack", "npush"])
 
 
-class AdaptDFA(object):
+class AdaptDFA:
     """Adapt Flow to the old DFA class expected by Interpreter"""
 
     def __init__(self, flow):
@@ -2543,7 +2543,7 @@ def _flatten_inst_regs(iterable):
                 yield x
 
 
-class AdaptCFA(object):
+class AdaptCFA:
     """Adapt Flow to the old CFA class expected by Interpreter"""
 
     def __init__(self, flow):
@@ -2585,7 +2585,7 @@ def dump(self):
         self._flow.cfgraph.dump()
 
 
-class AdaptCFBlock(object):
+class AdaptCFBlock:
     def __init__(self, blockinfo, offset):
         self.offset = offset
         self.body = tuple(i for i, _ in blockinfo.insts)
diff --git a/numba_cuda/numba/cuda/core/caching.py b/numba_cuda/numba/cuda/core/caching.py
index ed3ea1eb7..769259a85 100644
--- a/numba_cuda/numba/cuda/core/caching.py
+++ b/numba_cuda/numba/cuda/core/caching.py
@@ -89,7 +89,7 @@ def flush(self):
         pass
 
 
-class IndexDataCacheFile(object):
+class IndexDataCacheFile:
     """
     Implements the logic for the index file and data file used by a cache.
     """
@@ -407,7 +407,7 @@ def get_suitable_cache_subpath(cls, py_file):
         return "_".join([parentdir, hashed])
 
 
-class _SourceFileBackedLocatorMixin(object):
+class _SourceFileBackedLocatorMixin:
     """
     A cache locator mixin for functions which are backed by a well-known
     Python source file.
@@ -456,7 +456,7 @@ def get_cache_path(self):
         return self._cache_path
 
 
-class _SourceFileBackedLocatorMixin(object):
+class _SourceFileBackedLocatorMixin:
     """
     A cache locator mixin for functions which are backed by a well-known
     Python source file.
diff --git a/numba_cuda/numba/cuda/core/callconv.py b/numba_cuda/numba/cuda/core/callconv.py
index 2ff231934..c929c25b6 100644
--- a/numba_cuda/numba/cuda/core/callconv.py
+++ b/numba_cuda/numba/cuda/core/callconv.py
@@ -50,7 +50,7 @@ def _const_int(code):
 RETCODE_USEREXC = _const_int(FIRST_USEREXC)
 
 
-class BaseCallConv(object):
+class BaseCallConv:
     def __init__(self, context):
         self.context = context
 
@@ -293,7 +293,7 @@ def call_function(self, builder, callee, resty, argtys, args):
         return status, out
 
 
-class _MinimalCallHelper(object):
+class _MinimalCallHelper:
     """
     A call helper object for the "minimal" calling convention.
     User exceptions are represented as integer codes and stored in
@@ -341,7 +341,7 @@ def get_exception(self, exc_id):
             return exc, exc_args, locinfo
 
 
-class ErrorModel(object):
+class ErrorModel:
     def __init__(self, call_conv):
         self.call_conv = call_conv
 
diff --git a/numba_cuda/numba/cuda/core/compiler.py b/numba_cuda/numba/cuda/core/compiler.py
index f65e4b8f7..19523a33d 100644
--- a/numba_cuda/numba/cuda/core/compiler.py
+++ b/numba_cuda/numba/cuda/core/compiler.py
@@ -11,7 +11,7 @@
 from numba.cuda.core.targetconfig import ConfigStack
 
 
-class _CompileStatus(object):
+class _CompileStatus:
     """
     Describes the state of compilation. Used like a C record.
     """
@@ -72,7 +72,7 @@ def _make_subtarget(targetctx, flags):
     return targetctx.subtarget(**subtargetoptions)
 
 
-class CompilerBase(object):
+class CompilerBase:
     """
     Stores and manages states for the compiler
     """
diff --git a/numba_cuda/numba/cuda/core/compiler_lock.py b/numba_cuda/numba/cuda/core/compiler_lock.py
index f39cf5cce..d64da4e6a 100644
--- a/numba_cuda/numba/cuda/core/compiler_lock.py
+++ b/numba_cuda/numba/cuda/core/compiler_lock.py
@@ -15,7 +15,7 @@
 
 
 # Lock for the preventing multiple compiler execution
-class _CompilerLock(object):
+class _CompilerLock:
     def __init__(self):
         self._lock = threading.RLock()
 
@@ -46,7 +46,7 @@ def _acquire_compile_lock(*args, **kwargs):
 
 
 # Wrapper that coordinates both numba and numba-cuda compiler locks
-class _DualCompilerLock(object):
+class _DualCompilerLock:
     """Wrapper that coordinates both the numba-cuda and upstream numba compiler locks."""
 
     def __init__(self, cuda_lock, numba_lock):
diff --git a/numba_cuda/numba/cuda/core/compiler_machinery.py b/numba_cuda/numba/cuda/core/compiler_machinery.py
index 07b778c3c..bf866a471 100644
--- a/numba_cuda/numba/cuda/core/compiler_machinery.py
+++ b/numba_cuda/numba/cuda/core/compiler_machinery.py
@@ -23,7 +23,7 @@
 _termcolor = errors.termcolor()
 
 
-class SimpleTimer(object):
+class SimpleTimer:
     """
     A simple context managed timer
     """
@@ -111,7 +111,7 @@ def get_analysis(self, pass_name):
         return self._analysis[pass_name]
 
 
-class SSACompliantMixin(object):
+class SSACompliantMixin:
     """Mixin to indicate a pass is SSA form compliant. Nothing is asserted
     about this condition at present.
     """
@@ -137,7 +137,7 @@ class LoweringPass(CompilerPass):
     pass
 
 
-class AnalysisUsage(object):
+class AnalysisUsage:
     """This looks and behaves like LLVM's AnalysisUsage because its like that."""
 
     def __init__(self):
@@ -171,7 +171,7 @@ def debug_print(*args, **kwargs):
 pass_timings = namedtuple("pass_timings", "init run finalize")
 
 
-class PassManager(object):
+class PassManager:
     """
     The PassManager is a named instance of a particular compilation pipeline
     """
@@ -437,7 +437,7 @@ def walk(lkey, rmap):
 pass_info = namedtuple("pass_info", "pass_inst mutates_CFG analysis_only")
 
 
-class PassRegistry(object):
+class PassRegistry:
     """
     Pass registry singleton class.
     """
diff --git a/numba_cuda/numba/cuda/core/config.py b/numba_cuda/numba/cuda/core/config.py
index 33ba6c65e..16b440af0 100644
--- a/numba_cuda/numba/cuda/core/config.py
+++ b/numba_cuda/numba/cuda/core/config.py
@@ -124,7 +124,7 @@ def _process_opt_level(opt_level):
         return _OptLevel(opt_level)
 
 
-class _EnvReloader(object):
+class _EnvReloader:
     def __init__(self):
         self.reset()
 
diff --git a/numba_cuda/numba/cuda/core/consts.py b/numba_cuda/numba/cuda/core/consts.py
index 760e7e540..c322072df 100644
--- a/numba_cuda/numba/cuda/core/consts.py
+++ b/numba_cuda/numba/cuda/core/consts.py
@@ -8,7 +8,7 @@
 from numba.cuda.core import ir
 
 
-class ConstantInference(object):
+class ConstantInference:
     """
     A constant inference engine for a given interpreter.
     Inference inspects the IR to try and compute a compile-time constant for
diff --git a/numba_cuda/numba/cuda/core/controlflow.py b/numba_cuda/numba/cuda/core/controlflow.py
index a5bc153af..104c01596 100644
--- a/numba_cuda/numba/cuda/core/controlflow.py
+++ b/numba_cuda/numba/cuda/core/controlflow.py
@@ -23,7 +23,7 @@
     raise NotImplementedError(PYVERSION)
 
 
-class CFBlock(object):
+class CFBlock:
     def __init__(self, offset):
         self.offset = offset
         self.body = []
@@ -93,7 +93,7 @@ def _non_empty_items(self):
         return [(k, vs) for k, vs in sorted(self.items()) if vs]
 
 
-class CFGraph(object):
+class CFGraph:
     """
     Generic (almost) implementation of a Control Flow Graph.
     """
@@ -749,7 +749,7 @@ def __ne__(self, other):
         return not self.__eq__(other)
 
 
-class ControlFlowAnalysis(object):
+class ControlFlowAnalysis:
     """
     Attributes
     ----------
diff --git a/numba_cuda/numba/cuda/core/cuda_errors.py b/numba_cuda/numba/cuda/core/cuda_errors.py
index efca4e4ba..16d15c92a 100644
--- a/numba_cuda/numba/cuda/core/cuda_errors.py
+++ b/numba_cuda/numba/cuda/core/cuda_errors.py
@@ -57,11 +57,9 @@ def highlight(x):
                 return x
 
         if loc:
-            super(NumbaWarning, self).__init__(
-                highlight("%s\n%s\n" % (msg, loc.strformat()))
-            )
+            super().__init__(highlight("%s\n%s\n" % (msg, loc.strformat())))
         else:
-            super(NumbaWarning, self).__init__(highlight("%s" % (msg,)))
+            super().__init__(highlight("%s" % (msg,)))
 
 
 class NumbaPerformanceWarning(NumbaWarning):
@@ -249,7 +247,7 @@ def termcolor():
 else:
     from colorama import init, reinit, deinit, Fore, Style
 
-    class ColorShell(object):
+    class ColorShell:
         _has_initialized = False
 
         def __init__(self):
@@ -264,7 +262,7 @@ def __exit__(self, *exc_detail):
             Style.RESET_ALL  # noqa: B018
             deinit()
 
-    class reset_terminal(object):
+    class reset_terminal:
         def __init__(self):
             self._buf = bytearray(b"")
 
@@ -499,7 +497,7 @@ def wrapper(*args, **kwargs):
         return decorator
 
 
-class WarningsFixer(object):
+class WarningsFixer:
     """
     An object "fixing" warnings of a given category caught during
     certain phases.  The warnings can have their filename and lineno fixed,
@@ -575,7 +573,7 @@ def highlight(x):
             new_msg = "%s\n%s\n" % (msg, loc.strformat())
         else:
             new_msg = "%s" % (msg,)
-        super(NumbaError, self).__init__(highlight(new_msg))
+        super().__init__(highlight(new_msg))
 
     @property
     def contexts(self):
@@ -649,7 +647,7 @@ def __init__(self, name, loc=None):
             "The compiler failed to analyze the bytecode. "
             "Variable '%s' is not defined." % name
         )
-        super(NotDefinedError, self).__init__(msg, loc=loc)
+        super().__init__(msg, loc=loc)
 
 
 class VerificationError(IRError):
@@ -678,7 +676,7 @@ class LoweringError(NumbaError):
     """
 
     def __init__(self, msg, loc=None):
-        super(LoweringError, self).__init__(msg, loc=loc)
+        super().__init__(msg, loc=loc)
 
 
 class UnsupportedParforsError(NumbaError):
@@ -717,7 +715,7 @@ def __init__(self, value, attr, loc=None):
         else:
             msg = "Unknown attribute '{attr}' of type {type}"
             msg = msg.format(type=value, attr=attr)
-        super(UntypedAttributeError, self).__init__(msg, loc=loc)
+        super().__init__(msg, loc=loc)
 
 
 class ByteCodeSupportError(NumbaError):
@@ -726,7 +724,7 @@ class ByteCodeSupportError(NumbaError):
     """
 
     def __init__(self, msg, loc=None):
-        super(ByteCodeSupportError, self).__init__(msg, loc=loc)
+        super().__init__(msg, loc=loc)
 
 
 class CompilerError(NumbaError):
@@ -743,7 +741,7 @@ class ConstantInferenceError(NumbaError):
     """
 
     def __init__(self, value, loc=None):
-        super(ConstantInferenceError, self).__init__(value, loc=loc)
+        super().__init__(value, loc=loc)
 
 
 class InternalError(NumbaError):
@@ -752,7 +750,7 @@ class InternalError(NumbaError):
     """
 
     def __init__(self, exception):
-        super(InternalError, self).__init__(str(exception))
+        super().__init__(str(exception))
         self.old_exception = exception
 
 
@@ -805,7 +803,7 @@ def __init__(self, arg_indices, fold_arguments=None, loc=None):
             the ``args`` and ``kwargs``.
         loc : numba.ir.Loc or None
         """
-        super(ForceLiteralArg, self).__init__(
+        super().__init__(
             "Pseudo-exception to force literal arguments in the dispatcher",
             loc=loc,
         )
diff --git a/numba_cuda/numba/cuda/core/funcdesc.py b/numba_cuda/numba/cuda/core/funcdesc.py
index bcb238d00..f6ed3638d 100644
--- a/numba_cuda/numba/cuda/core/funcdesc.py
+++ b/numba_cuda/numba/cuda/core/funcdesc.py
@@ -25,7 +25,7 @@ def qualifying_prefix(modname, qualname):
     return "{}.{}".format(modname, qualname) if modname else qualname
 
 
-class FunctionDescriptor(object):
+class FunctionDescriptor:
     """
     Base class for function descriptors: an object used to carry
     useful metadata about a natively callable function.
@@ -314,7 +314,7 @@ def __init__(self, name, restype, argtypes):
         def mangler(a, x, abi_tags, uid=None):
             return a
 
-        super(ExternalFunctionDescriptor, self).__init__(
+        super().__init__(
             native=True,
             modname=None,
             qualname=name,
diff --git a/numba_cuda/numba/cuda/core/generators.py b/numba_cuda/numba/cuda/core/generators.py
index f296d6b04..63d68dbed 100644
--- a/numba_cuda/numba/cuda/core/generators.py
+++ b/numba_cuda/numba/cuda/core/generators.py
@@ -61,7 +61,7 @@ def llvm_finalizer_name(self):
         return "finalize_" + self.mangled_name
 
 
-class BaseGeneratorLower(object):
+class BaseGeneratorLower:
     """
     Base support class for lowering generators.
     """
@@ -323,7 +323,7 @@ def lower_finalize_func_body(self, builder, genptr):
         builder.ret_void()
 
 
-class LowerYield(object):
+class LowerYield:
     """
     Support class for lowering a particular yield point.
     """
diff --git a/numba_cuda/numba/cuda/core/imputils.py b/numba_cuda/numba/cuda/core/imputils.py
index 2881c2771..8cb07adb9 100644
--- a/numba_cuda/numba/cuda/core/imputils.py
+++ b/numba_cuda/numba/cuda/core/imputils.py
@@ -14,7 +14,7 @@
 from numba.cuda.typing.templates import BaseRegistryLoader
 
 
-class Registry(object):
+class Registry:
     """
     A registry of function and attribute implementations.
     """
@@ -275,7 +275,7 @@ def iternext_wrapper(context, builder, sig, args, result):
     return wrapper
 
 
-class _IternextResult(object):
+class _IternextResult:
     """
     A result wrapper for iteration, passed by iternext_impl() into the
     wrapped function.
diff --git a/numba_cuda/numba/cuda/core/inline_closurecall.py b/numba_cuda/numba/cuda/core/inline_closurecall.py
index a6e952fe2..3f6bfa6f8 100644
--- a/numba_cuda/numba/cuda/core/inline_closurecall.py
+++ b/numba_cuda/numba/cuda/core/inline_closurecall.py
@@ -74,7 +74,7 @@ def _created_inlined_var_name(function_name, var_name):
     return new_name
 
 
-class InlineClosureCallPass(object):
+class InlineClosureCallPass:
     """InlineClosureCallPass class looks for direct calls to locally defined
     closures, and inlines the body of the closure function to the call site.
     """
@@ -259,7 +259,7 @@ def check_reduce_func(func_ir, func_var):
     return reduce_func
 
 
-class InlineWorker(object):
+class InlineWorker:
     """A worker class for inlining, this is a more advanced version of
     `inline_closure_call` in that it permits inlining from function type, Numba
     IR and code object. It also, runs the entire untyped compiler pipeline on
@@ -1555,7 +1555,7 @@ class RewriteArrayOfConsts(rewrites.Rewrite):
 
     def __init__(self, state, *args, **kws):
         self.typingctx = state.typingctx
-        super(RewriteArrayOfConsts, self).__init__(*args, **kws)
+        super().__init__(*args, **kws)
 
     def match(self, func_ir, block, typemap, calltypes):
         if len(calltypes) == 0:
@@ -1692,7 +1692,7 @@ def inline_array(array_var, expr, stmts, list_vars, dels):
         stmts.extend(dels)
         return True
 
-    class State(object):
+    class State:
         """
         This class is used to hold the state in the following loop so as to make
         it easy to reset the state of the variables tracking the various
diff --git a/numba_cuda/numba/cuda/core/interpreter.py b/numba_cuda/numba/cuda/core/interpreter.py
index 0a7634715..2db457e7a 100644
--- a/numba_cuda/numba/cuda/core/interpreter.py
+++ b/numba_cuda/numba/cuda/core/interpreter.py
@@ -36,7 +36,7 @@
     raise NotImplementedError(PYVERSION)
 
 
-class _UNKNOWN_VALUE(object):
+class _UNKNOWN_VALUE:
     """Represents an unknown value, this is for ease of debugging purposes only."""
 
     def __init__(self, varname):
@@ -49,7 +49,7 @@ def __repr__(self):
 _logger = logging.getLogger(__name__)
 
 
-class Assigner(object):
+class Assigner:
     """
     This object keeps track of potential assignment simplifications
     inside a code block.
@@ -1347,7 +1347,7 @@ def _build_new_build_map(func_ir, name, old_body, old_lineno, new_items):
     )
 
 
-class Interpreter(object):
+class Interpreter:
     """A bytecode interpreter that builds up the IR."""
 
     _DEBUG_PRINT = False
diff --git a/numba_cuda/numba/cuda/core/ir.py b/numba_cuda/numba/cuda/core/ir.py
index 901473aba..caa89fd20 100644
--- a/numba_cuda/numba/cuda/core/ir.py
+++ b/numba_cuda/numba/cuda/core/ir.py
@@ -33,7 +33,7 @@
 _termcolor = errors.termcolor()
 
 
-class Loc(object):
+class Loc:
     """Source location"""
 
     _defmatcher = re.compile(r"def\s+(\w+)")
@@ -226,7 +226,7 @@ def short(self):
 
 
 @total_ordering
-class SlotEqualityCheckMixin(object):
+class SlotEqualityCheckMixin:
     # some ir nodes are __dict__ free using __slots__ instead, this mixin
     # should not trigger the unintended creation of __dict__.
     __slots__ = tuple()
@@ -248,7 +248,7 @@ def __hash__(self):
 
 
 @total_ordering
-class EqualityCheckMixin(object):
+class EqualityCheckMixin:
     """Mixin for basic equality checking"""
 
     def __eq__(self, other):
@@ -274,7 +274,7 @@ def __hash__(self):
         return id(self)
 
 
-class VarMap(object):
+class VarMap:
     def __init__(self):
         self._con = {}
 
@@ -315,7 +315,7 @@ def __ne__(self, other):
         return not self.__eq__(other)
 
 
-class AbstractRHS(object):
+class AbstractRHS:
     """Abstract base class for anything that can be the RHS of an assignment.
     This class **does not** define any methods.
     """
@@ -1475,7 +1475,7 @@ def __repr__(self):
         return "With(entry=%s, exit=%s)" % args
 
 
-class FunctionIR(object):
+class FunctionIR:
     def __init__(
         self,
         blocks,
diff --git a/numba_cuda/numba/cuda/core/ir_utils.py b/numba_cuda/numba/cuda/core/ir_utils.py
index 6408bfa7a..14cd4bed0 100644
--- a/numba_cuda/numba/cuda/core/ir_utils.py
+++ b/numba_cuda/numba/cuda/core/ir_utils.py
@@ -1992,7 +1992,7 @@ def get_ir_of_code(glbls, fcode):
     # we need to run the before inference rewrite pass to normalize the IR
     # XXX: check rewrite pass flag?
     # for example, Raise nodes need to become StaticRaise before type inference
-    class DummyPipeline(object):
+    class DummyPipeline:
         def __init__(self, f_ir):
             self.state = StateDict()
             self.state.typingctx = None
diff --git a/numba_cuda/numba/cuda/core/postproc.py b/numba_cuda/numba/cuda/core/postproc.py
index 61b51c7a7..78a626706 100644
--- a/numba_cuda/numba/cuda/core/postproc.py
+++ b/numba_cuda/numba/cuda/core/postproc.py
@@ -5,7 +5,7 @@
 from numba.cuda.core import ir_utils, transforms
 
 
-class YieldPoint(object):
+class YieldPoint:
     def __init__(self, block, inst):
         assert isinstance(block, ir.block_types)
         assert isinstance(inst, ir.yield_types)
@@ -15,7 +15,7 @@ def __init__(self, block, inst):
         self.weak_live_vars = None
 
 
-class GeneratorInfo(object):
+class GeneratorInfo:
     def __init__(self):
         # { index: YieldPoint }
         self.yield_points = {}
@@ -29,7 +29,7 @@ def get_yield_points(self):
         return self.yield_points.values()
 
 
-class VariableLifetime(object):
+class VariableLifetime:
     """
     For lazily building information of variable lifetime
     """
@@ -63,7 +63,7 @@ def deadmaps(self):
 ir_extension_insert_dels = {}
 
 
-class PostProcessor(object):
+class PostProcessor:
     """
     A post-processor for Numba IR.
     """
diff --git a/numba_cuda/numba/cuda/core/pythonapi.py b/numba_cuda/numba/cuda/core/pythonapi.py
index 592e25780..749b3651d 100644
--- a/numba_cuda/numba/cuda/core/pythonapi.py
+++ b/numba_cuda/numba/cuda/core/pythonapi.py
@@ -23,7 +23,7 @@
     PY_UNICODE_WCHAR_KIND = _helperlib.py_unicode_wchar_kind
 
 
-class _Registry(object):
+class _Registry:
     def __init__(self):
         self.functions = {}
 
@@ -106,7 +106,7 @@ def reflect(self, typ, val):
         return self.pyapi.reflect_native_value(typ, val, self.env_manager)
 
 
-class NativeValue(object):
+class NativeValue:
     """
     Encapsulate the result of converting a Python object to a native value,
     recording whether the conversion was successful and how to cleanup.
@@ -118,7 +118,7 @@ def __init__(self, value, is_error=None, cleanup=None):
         self.cleanup = cleanup
 
 
-class EnvironmentManager(object):
+class EnvironmentManager:
     def __init__(self, pyapi, env, env_body, env_ptr):
         assert isinstance(env, lowering.Environment)
         self.pyapi = pyapi
@@ -174,7 +174,7 @@ def read_const(self, index):
 _IteratorLoop = namedtuple("_IteratorLoop", ("value", "do_break"))
 
 
-class PythonAPI(object):
+class PythonAPI:
     """
     Code generation facilities to call into the CPython C API (and related
     helpers).
diff --git a/numba_cuda/numba/cuda/core/registry.py b/numba_cuda/numba/cuda/core/registry.py
index c8beef953..4894f1f60 100644
--- a/numba_cuda/numba/cuda/core/registry.py
+++ b/numba_cuda/numba/cuda/core/registry.py
@@ -22,13 +22,13 @@ def __init__(self, *args, **kws):
         self.key_type = kws.pop("key_type", None)
         self.value_type = kws.pop("value_type", None)
         self._type_check = self.key_type or self.value_type
-        super(DelayedRegistry, self).__init__(*args, **kws)
+        super().__init__(*args, **kws)
 
     def __getitem__(self, item):
         if item in self.ondemand:
             self[item] = self.ondemand[item]()
             del self.ondemand[item]
-        return super(DelayedRegistry, self).__getitem__(item)
+        return super().__getitem__(item)
 
     def __setitem__(self, key, value):
         if self._type_check:
@@ -43,4 +43,4 @@ def check(x, ty_x):
                 check(key, self.key_type)
             if self.value_type is not None:
                 check(value, self.value_type)
-        return super(DelayedRegistry, self).__setitem__(key, value)
+        return super().__setitem__(key, value)
diff --git a/numba_cuda/numba/cuda/core/rewrites/registry.py b/numba_cuda/numba/cuda/core/rewrites/registry.py
index 2ed67dcaa..dbcb94c44 100644
--- a/numba_cuda/numba/cuda/core/rewrites/registry.py
+++ b/numba_cuda/numba/cuda/core/rewrites/registry.py
@@ -6,7 +6,7 @@
 from numba.cuda import config
 
 
-class Rewrite(object):
+class Rewrite:
     """Defines the abstract base class for Numba rewrites."""
 
     def __init__(self, state=None):
@@ -26,7 +26,7 @@ def apply(self):
         raise NotImplementedError("Abstract Rewrite.apply() called!")
 
 
-class RewriteRegistry(object):
+class RewriteRegistry:
     """Defines a registry for Numba rewrites."""
 
     _kinds = frozenset(["before-inference", "after-inference"])
diff --git a/numba_cuda/numba/cuda/core/typeinfer.py b/numba_cuda/numba/cuda/core/typeinfer.py
index ccc21e050..08159cbf2 100644
--- a/numba_cuda/numba/cuda/core/typeinfer.py
+++ b/numba_cuda/numba/cuda/core/typeinfer.py
@@ -63,7 +63,7 @@ class NOTSET:
 _termcolor = termcolor()
 
 
-class TypeVar(object):
+class TypeVar:
     def __init__(self, context, var):
         self.context = context
         self.var = var
@@ -163,7 +163,7 @@ def __len__(self):
         return 1 if self.type is not None else 0
 
 
-class ConstraintNetwork(object):
+class ConstraintNetwork:
     """
     TODO: It is possible to optimize constraint propagation to consider only
           dirty type variables.
@@ -204,7 +204,7 @@ def propagate(self, typeinfer):
         return errors
 
 
-class Propagate(object):
+class Propagate:
     """
     A simple constraint for direct propagation of types for assignments.
     """
@@ -230,7 +230,7 @@ def refine(self, typeinfer, target_type):
         )
 
 
-class ArgConstraint(object):
+class ArgConstraint:
     def __init__(self, dst, src, loc):
         self.dst = dst
         self.src = src
@@ -252,7 +252,7 @@ def __call__(self, typeinfer):
             typeinfer.add_type(self.dst, ty, loc=self.loc)
 
 
-class BuildTupleConstraint(object):
+class BuildTupleConstraint:
     def __init__(self, target, items, loc):
         self.target = target
         self.items = items
@@ -272,7 +272,7 @@ def __call__(self, typeinfer):
                 typeinfer.add_type(self.target, tup, loc=self.loc)
 
 
-class _BuildContainerConstraint(object):
+class _BuildContainerConstraint:
     def __init__(self, target, items, loc):
         self.target = target
         self.items = items
@@ -339,7 +339,7 @@ class BuildSetConstraint(_BuildContainerConstraint):
     container_type = types.Set
 
 
-class BuildMapConstraint(object):
+class BuildMapConstraint:
     def __init__(self, target, items, special_value, value_indexes, loc):
         self.target = target
         self.items = items
@@ -417,7 +417,7 @@ def check(other):
                     )
 
 
-class ExhaustIterConstraint(object):
+class ExhaustIterConstraint:
     def __init__(self, target, count, iterator, loc):
         self.target = target
         self.count = count
@@ -454,7 +454,7 @@ def __call__(self, typeinfer):
                     )
 
 
-class PairFirstConstraint(object):
+class PairFirstConstraint:
     def __init__(self, target, pair, loc):
         self.target = target
         self.pair = pair
@@ -474,7 +474,7 @@ def __call__(self, typeinfer):
                 typeinfer.add_type(self.target, tp.first_type, loc=self.loc)
 
 
-class PairSecondConstraint(object):
+class PairSecondConstraint:
     def __init__(self, target, pair, loc):
         self.target = target
         self.pair = pair
@@ -491,7 +491,7 @@ def __call__(self, typeinfer):
                 typeinfer.add_type(self.target, tp.second_type, loc=self.loc)
 
 
-class StaticGetItemConstraint(object):
+class StaticGetItemConstraint:
     def __init__(self, target, value, index, index_var, loc):
         self.target = target
         self.value = value
@@ -526,7 +526,7 @@ def get_call_signature(self):
         return self.fallback and self.fallback.get_call_signature()
 
 
-class TypedGetItemConstraint(object):
+class TypedGetItemConstraint:
     def __init__(self, target, value, dtype, index, loc):
         self.target = target
         self.value = value
@@ -593,7 +593,7 @@ def _is_array_not_precise(arrty):
     return isinstance(arrty, types.Array) and not arrty.is_precise()
 
 
-class CallConstraint(object):
+class CallConstraint:
     """Constraint for calling functions.
     Perform case analysis foreach combinations of argument types.
     """
@@ -755,7 +755,7 @@ def __call__(self, typeinfer):
             self.resolve(typeinfer, typeinfer.typevars, fnty=fnty)
 
 
-class GetAttrConstraint(object):
+class GetAttrConstraint:
     def __init__(self, target, attr, value, loc, inst):
         self.target = target
         self.attr = attr
@@ -793,7 +793,7 @@ def __repr__(self):
         )
 
 
-class SetItemRefinement(object):
+class SetItemRefinement:
     """A mixin class to provide the common refinement logic in setitem
     and static setitem.
     """
@@ -892,7 +892,7 @@ def get_call_signature(self):
         return self.signature
 
 
-class DelItemConstraint(object):
+class DelItemConstraint:
     def __init__(self, target, index, loc):
         self.target = target
         self.index = index
@@ -920,7 +920,7 @@ def get_call_signature(self):
         return self.signature
 
 
-class SetAttrConstraint(object):
+class SetAttrConstraint:
     def __init__(self, target, attr, value, loc):
         self.target = target
         self.attr = attr
@@ -951,7 +951,7 @@ def get_call_signature(self):
         return self.signature
 
 
-class PrintConstraint(object):
+class PrintConstraint:
     def __init__(self, args, vararg, loc):
         self.args = args
         self.vararg = vararg
@@ -982,14 +982,14 @@ def set_context(self, context):
     def __getitem__(self, name):
         if name not in self:
             self[name] = TypeVar(self.context, name)
-        return super(TypeVarMap, self).__getitem__(name)
+        return super().__getitem__(name)
 
     def __setitem__(self, name, value):
         assert isinstance(name, str)
         if name in self:
             raise KeyError("Cannot redefine typevar %s" % name)
         else:
-            super(TypeVarMap, self).__setitem__(name, value)
+            super().__setitem__(name, value)
 
 
 # A temporary mapping of {function name: dispatcher object}
@@ -1023,7 +1023,7 @@ def register_dispatcher(disp):
 typeinfer_extensions = {}
 
 
-class TypeInferer(object):
+class TypeInferer:
     """
     Operates on block that shares the same ir.Scope.
     """
@@ -1931,7 +1931,7 @@ def typeof_intrinsic_call(self, inst, target, func, *args):
         self.calls.append((inst.value, constraint))
 
 
-class NullDebug(object):
+class NullDebug:
     def propagate_started(self):
         pass
 
@@ -1942,7 +1942,7 @@ def unify_finished(self, typdict, retty, fntys):
         pass
 
 
-class TypeInferDebug(object):
+class TypeInferDebug:
     def __init__(self, typeinfer):
         self.typeinfer = typeinfer
 
diff --git a/numba_cuda/numba/cuda/cpython/listobj.py b/numba_cuda/numba/cuda/cpython/listobj.py
index ee059a8e7..68be31663 100644
--- a/numba_cuda/numba/cuda/cpython/listobj.py
+++ b/numba_cuda/numba/cuda/cpython/listobj.py
@@ -47,7 +47,7 @@ def get_itemsize(context, list_type):
     return context.get_abi_sizeof(llty)
 
 
-class _ListPayloadMixin(object):
+class _ListPayloadMixin:
     @property
     def size(self):
         return self._payload.size
diff --git a/numba_cuda/numba/cuda/cpython/unicode.py b/numba_cuda/numba/cuda/cpython/unicode.py
index d6e194a49..91bfc539a 100644
--- a/numba_cuda/numba/cuda/cpython/unicode.py
+++ b/numba_cuda/numba/cuda/cpython/unicode.py
@@ -124,7 +124,7 @@ def __init__(self, dmm, fe_type):
             ("index", types.EphemeralPointer(types.uintp)),
             ("data", fe_type.data),
         ]
-        super(UnicodeIteratorModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 # CAST
diff --git a/numba_cuda/numba/cuda/cudadrv/devicearray.py b/numba_cuda/numba/cuda/cudadrv/devicearray.py
index 2f82b1414..7a313f4e0 100644
--- a/numba_cuda/numba/cuda/cudadrv/devicearray.py
+++ b/numba_cuda/numba/cuda/cudadrv/devicearray.py
@@ -429,9 +429,7 @@ class DeviceRecord(DeviceNDArrayBase):
     def __init__(self, dtype, stream=0, gpu_data=None):
         shape = ()
         strides = ()
-        super(DeviceRecord, self).__init__(
-            shape, strides, dtype, stream, gpu_data
-        )
+        super().__init__(shape, strides, dtype, stream, gpu_data)
 
     @property
     def flags(self):
@@ -784,7 +782,7 @@ def _do_setitem(self, key, value, stream=0):
             stream.synchronize()
 
 
-class IpcArrayHandle(object):
+class IpcArrayHandle:
     """
     An IPC array handle that can be serialized and transfer to another process
     in the same machine for share a GPU allocation.
diff --git a/numba_cuda/numba/cuda/cudadrv/devices.py b/numba_cuda/numba/cuda/cudadrv/devices.py
index 514191384..0795cc6e5 100644
--- a/numba_cuda/numba/cuda/cudadrv/devices.py
+++ b/numba_cuda/numba/cuda/cudadrv/devices.py
@@ -56,7 +56,7 @@ def current(self):
                 return self[devnum]
 
 
-class _DeviceContextManager(object):
+class _DeviceContextManager:
     """
     Provides a context manager for executing in the context of the chosen
     device. The normal use of instances of this type is from
@@ -88,7 +88,7 @@ def __str__(self):
         return f"<Managed Device {self.id}>"
 
 
-class _Runtime(object):
+class _Runtime:
     """Emulate the CUDA runtime context management.
 
     It owns all Devices and Contexts.
diff --git a/numba_cuda/numba/cuda/cudadrv/driver.py b/numba_cuda/numba/cuda/cudadrv/driver.py
index 2dd4cb276..cc84a385e 100644
--- a/numba_cuda/numba/cuda/cudadrv/driver.py
+++ b/numba_cuda/numba/cuda/cudadrv/driver.py
@@ -137,7 +137,7 @@ class CudaAPIError(CudaDriverError):
     def __init__(self, code, msg):
         self.code = code
         self.msg = msg
-        super(CudaAPIError, self).__init__(code, msg)
+        super().__init__(code, msg)
 
     def __str__(self):
         return "[%s] %s" % (self.code, self.msg)
@@ -233,7 +233,7 @@ def _getpid():
 ERROR_MAP = _build_reverse_error_map()
 
 
-class Driver(object):
+class Driver:
     """
     Driver API functions are lazily bound.
     """
@@ -453,7 +453,7 @@ def get_version(self):
         return (major, minor)
 
 
-class _ActiveContext(object):
+class _ActiveContext:
     """An contextmanager object to cache active context to reduce dependency
     on querying the CUDA driver API.
 
@@ -1006,7 +1006,7 @@ def __str__(self):
 _SizeNotSet = _SizeNotSet()
 
 
-class _PendingDeallocs(object):
+class _PendingDeallocs:
     """
     Pending deallocations of a context (or device since we are using the primary
     context). The capacity defaults to being unset (_SizeNotSet) but can be
@@ -1092,7 +1092,7 @@ def __len__(self):
 """
 
 
-class Context(object):
+class Context:
     """
     This object wraps a CUDA Context resource.
 
@@ -1435,7 +1435,7 @@ def module_unload(handle):
     return core
 
 
-class _CudaIpcImpl(object):
+class _CudaIpcImpl:
     """Implementation of GPU IPC using CUDA driver API.
     This requires the devices to be peer accessible.
     """
@@ -1472,7 +1472,7 @@ def close(self):
         self._opened_mem = None
 
 
-class _StagedIpcImpl(object):
+class _StagedIpcImpl:
     """Implementation of GPU IPC using custom staging logic to workaround
     CUDA IPC limitation on peer accessibility between devices.
     """
@@ -1512,7 +1512,7 @@ def close(self):
         pass
 
 
-class IpcHandle(object):
+class IpcHandle:
     """
     CUDA IPC handle. Serialization of the CUDA IPC handle object is implemented
     here.
@@ -1633,7 +1633,7 @@ def _rebuild(cls, handle_ary, size, source_info, offset):
         )
 
 
-class MemoryPointer(object):
+class MemoryPointer:
     """A memory pointer that owns a buffer, with an optional finalizer. Memory
     pointers provide reference counting, and instances are initialized with a
     reference count of 1.
@@ -1756,7 +1756,7 @@ class AutoFreePointer(MemoryPointer):
     """
 
     def __init__(self, *args, **kwargs):
-        super(AutoFreePointer, self).__init__(*args, **kwargs)
+        super().__init__(*args, **kwargs)
         # Releease the self reference to the buffer, so that the finalizer
         # is invoked if all the derived pointers are gone.
         self.refct -= 1
@@ -1792,9 +1792,7 @@ def __init__(self, context, pointer, size, owner=None, finalizer=None):
         self._bufptr_ = self.host_pointer
 
         self.device_pointer = devptr
-        super(MappedMemory, self).__init__(
-            context, devptr, size, finalizer=finalizer
-        )
+        super().__init__(context, devptr, size, finalizer=finalizer)
         self.handle = self.host_pointer
 
         # For buffer interface
@@ -1877,7 +1875,7 @@ def own(self):
         return ManagedOwnedPointer(weakref.proxy(self))
 
 
-class OwnedPointer(object):
+class OwnedPointer:
     def __init__(self, memptr, view=None):
         self._mem = memptr
 
diff --git a/numba_cuda/numba/cuda/cudadrv/dummyarray.py b/numba_cuda/numba/cuda/cudadrv/dummyarray.py
index 165459545..73a2fc6b1 100644
--- a/numba_cuda/numba/cuda/cudadrv/dummyarray.py
+++ b/numba_cuda/numba/cuda/cudadrv/dummyarray.py
@@ -120,7 +120,7 @@ def attempt_nocopy_reshape(
     return 1
 
 
-class Dim(object):
+class Dim:
     """A single dimension of the array
 
     Attributes
@@ -208,7 +208,7 @@ def compute_index(indices, dims):
     return sum(d.get_offset(i) for i, d in zip(indices, dims))
 
 
-class Element(object):
+class Element:
     is_array = False
 
     def __init__(self, extent):
@@ -218,7 +218,7 @@ def iter_contiguous_extent(self):
         yield self.extent
 
 
-class Array(object):
+class Array:
     """A dummy numpy array-like object.  Consider it an array without the
     actual data, but offset from the base data pointer.
 
diff --git a/numba_cuda/numba/cuda/cudadrv/nvvm.py b/numba_cuda/numba/cuda/cudadrv/nvvm.py
index 2d0ad4c6c..d6ddc27ca 100644
--- a/numba_cuda/numba/cuda/cudadrv/nvvm.py
+++ b/numba_cuda/numba/cuda/cudadrv/nvvm.py
@@ -72,7 +72,7 @@ def is_available():
 _nvvm_lock = threading.Lock()
 
 
-class NVVM(object):
+class NVVM:
     """Process-wide singleton."""
 
     _PROTOTYPES = {
@@ -274,7 +274,7 @@ def check_error(self, error, msg, exit=False):
                 raise exc
 
 
-class CompilationUnit(object):
+class CompilationUnit:
     """
     A CompilationUnit is a set of LLVM modules that are compiled to PTX or
     LTO-IR with NVVM.
@@ -420,7 +420,7 @@ def get_log(self):
 """
 
 
-class LibDevice(object):
+class LibDevice:
     _cache_ = None
 
     def __init__(self):
diff --git a/numba_cuda/numba/cuda/datamodel/cuda_manager.py b/numba_cuda/numba/cuda/datamodel/cuda_manager.py
index 90d0cebec..7c1d724dc 100644
--- a/numba_cuda/numba/cuda/datamodel/cuda_manager.py
+++ b/numba_cuda/numba/cuda/datamodel/cuda_manager.py
@@ -7,7 +7,7 @@
 from numba.cuda import types
 
 
-class DataModelManager(object):
+class DataModelManager:
     """Manages mapping of FE types to their corresponding data model"""
 
     def __init__(self, handlers=None):
diff --git a/numba_cuda/numba/cuda/datamodel/cuda_models.py b/numba_cuda/numba/cuda/datamodel/cuda_models.py
index 4b0dab47c..6108ae8d0 100644
--- a/numba_cuda/numba/cuda/datamodel/cuda_models.py
+++ b/numba_cuda/numba/cuda/datamodel/cuda_models.py
@@ -12,7 +12,7 @@
 from numba.cuda.np import numpy_support
 
 
-class DataModel(object):
+class DataModel:
     """
     DataModel describe how a FE type is represented in the LLVM IR at
     different contexts.
@@ -232,7 +232,7 @@ class PrimitiveModel(DataModel):
     """
 
     def __init__(self, dmm, fe_type, be_type):
-        super(PrimitiveModel, self).__init__(dmm, fe_type)
+        super().__init__(dmm, fe_type)
         self.be_type = be_type
 
     def get_value_type(self):
@@ -301,7 +301,7 @@ class EnumModel(ProxyModel):
     """
 
     def __init__(self, dmm, fe_type):
-        super(EnumModel, self).__init__(dmm, fe_type)
+        super().__init__(dmm, fe_type)
         self._proxied_model = dmm.lookup(fe_type.dtype)
 
 
@@ -342,7 +342,7 @@ class OpaqueModel(PrimitiveModel):
 
     def __init__(self, dmm, fe_type):
         be_type = self._ptr_type
-        super(OpaqueModel, self).__init__(dmm, fe_type, be_type)
+        super().__init__(dmm, fe_type, be_type)
 
 
 @register_default(types.MemInfoPointer)
@@ -362,7 +362,7 @@ def get_nrt_meminfo(self, builder, value):
 class IntegerModel(PrimitiveModel):
     def __init__(self, dmm, fe_type):
         be_type = ir.IntType(fe_type.bitwidth)
-        super(IntegerModel, self).__init__(dmm, fe_type, be_type)
+        super().__init__(dmm, fe_type, be_type)
 
 
 @register_default(types.Float)
@@ -374,7 +374,7 @@ def __init__(self, dmm, fe_type):
             be_type = ir.DoubleType()
         else:
             raise NotImplementedError(fe_type)
-        super(FloatModel, self).__init__(dmm, fe_type, be_type)
+        super().__init__(dmm, fe_type, be_type)
 
 
 @register_default(types.CPointer)
@@ -383,7 +383,7 @@ def __init__(self, dmm, fe_type):
         self._pointee_model = dmm.lookup(fe_type.dtype)
         self._pointee_be_type = self._pointee_model.get_data_type()
         be_type = self._pointee_be_type.as_pointer()
-        super(PointerModel, self).__init__(dmm, fe_type, be_type)
+        super().__init__(dmm, fe_type, be_type)
 
 
 @register_default(types.EphemeralPointer)
@@ -405,7 +405,7 @@ def load_from_data_pointer(self, builder, ptr, align=None):
 @register_default(types.EphemeralArray)
 class EphemeralArrayModel(PointerModel):
     def __init__(self, dmm, fe_type):
-        super(EphemeralArrayModel, self).__init__(dmm, fe_type)
+        super().__init__(dmm, fe_type)
         self._data_type = ir.ArrayType(
             self._pointee_be_type, self._fe_type.count
         )
@@ -436,7 +436,7 @@ def __init__(self, dmm, fe_type):
         retty = dmm.lookup(sig.return_type).get_value_type()
         args = [dmm.lookup(t).get_value_type() for t in sig.args]
         be_type = ir.PointerType(ir.FunctionType(retty, args))
-        super(ExternalFuncPointerModel, self).__init__(dmm, fe_type, be_type)
+        super().__init__(dmm, fe_type, be_type)
 
 
 @register_default(types.UniTuple)
@@ -444,7 +444,7 @@ def __init__(self, dmm, fe_type):
 @register_default(types.StarArgUniTuple)
 class UniTupleModel(DataModel):
     def __init__(self, dmm, fe_type):
-        super(UniTupleModel, self).__init__(dmm, fe_type)
+        super().__init__(dmm, fe_type)
         self._elem_model = dmm.lookup(fe_type.dtype)
         self._count = len(fe_type)
         self._value_type = ir.ArrayType(
@@ -529,7 +529,7 @@ class StructModel(CompositeModel):
     _data_type = None
 
     def __init__(self, dmm, fe_type, members):
-        super(StructModel, self).__init__(dmm, fe_type)
+        super().__init__(dmm, fe_type)
         if members:
             self._fields, self._members = zip(*members)
         else:
@@ -750,7 +750,7 @@ def __init__(self, dmm, fe_type):
             ("real", fe_type.underlying_float),
             ("imag", fe_type.underlying_float),
         ]
-        super(ComplexModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.LiteralList)
@@ -761,7 +761,7 @@ def __init__(self, dmm, fe_type):
 class TupleModel(StructModel):
     def __init__(self, dmm, fe_type):
         members = [("f" + str(i), t) for i, t in enumerate(fe_type)]
-        super(TupleModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.UnionType)
@@ -772,7 +772,7 @@ def __init__(self, dmm, fe_type):
             # XXX: it should really be a MemInfoPointer(types.voidptr)
             ("payload", types.Tuple.from_types(fe_type.types)),
         ]
-        super(UnionModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.Pair)
@@ -782,7 +782,7 @@ def __init__(self, dmm, fe_type):
             ("first", fe_type.first_type),
             ("second", fe_type.second_type),
         ]
-        super(PairModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.ListPayload)
@@ -799,7 +799,7 @@ def __init__(self, dmm, fe_type):
             # Actually an inlined var-sized array
             ("data", fe_type.container.dtype),
         ]
-        super(ListPayloadModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.List)
@@ -812,7 +812,7 @@ def __init__(self, dmm, fe_type):
             # This member is only used only for reflected lists
             ("parent", types.pyobject),
         ]
-        super(ListModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.ListIter)
@@ -825,7 +825,7 @@ def __init__(self, dmm, fe_type):
             ("meminfo", types.MemInfoPointer(payload_type)),
             ("index", types.EphemeralPointer(types.intp)),
         ]
-        super(ListIterModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.SetEntry)
@@ -837,7 +837,7 @@ def __init__(self, dmm, fe_type):
             ("hash", types.intp),
             ("key", dtype),
         ]
-        super(SetEntryModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.SetPayload)
@@ -858,7 +858,7 @@ def __init__(self, dmm, fe_type):
             # Actually an inlined var-sized array
             ("entries", entry_type),
         ]
-        super(SetPayloadModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.Set)
@@ -871,7 +871,7 @@ def __init__(self, dmm, fe_type):
             # This member is only used only for reflected sets
             ("parent", types.pyobject),
         ]
-        super(SetModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.SetIter)
@@ -885,7 +885,7 @@ def __init__(self, dmm, fe_type):
             # The index into the entries table
             ("index", types.EphemeralPointer(types.intp)),
         ]
-        super(SetIterModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.Array)
@@ -906,7 +906,7 @@ def __init__(self, dmm, fe_type):
             ("shape", types.UniTuple(types.intp, ndim)),
             ("strides", types.UniTuple(types.intp, ndim)),
         ]
-        super(ArrayModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.ArrayFlags)
@@ -915,14 +915,14 @@ def __init__(self, dmm, fe_type):
         members = [
             ("parent", fe_type.array_type),
         ]
-        super(ArrayFlagsModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.NestedArray)
 class NestedArrayModel(ArrayModel):
     def __init__(self, dmm, fe_type):
         self._be_type = dmm.lookup(fe_type.dtype).get_data_type()
-        super(NestedArrayModel, self).__init__(dmm, fe_type)
+        super().__init__(dmm, fe_type)
 
     def as_storage_type(self):
         """Return the LLVM type representation for the storage of
@@ -940,7 +940,7 @@ def __init__(self, dmm, fe_type):
             ("valid", types.boolean),
         ]
         self._value_model = dmm.lookup(fe_type.type)
-        super(OptionalModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
     def get_return_type(self):
         return self._value_model.get_return_type()
@@ -969,7 +969,7 @@ def get_valid(value):
 @register_default(types.Record)
 class RecordModel(CompositeModel):
     def __init__(self, dmm, fe_type):
-        super(RecordModel, self).__init__(dmm, fe_type)
+        super().__init__(dmm, fe_type)
         self._models = [self._dmm.lookup(t) for _, t in fe_type.members]
         self._be_type = ir.ArrayType(ir.IntType(8), fe_type.size)
         self._be_ptr_type = self._be_type.as_pointer()
@@ -1012,7 +1012,7 @@ def load_from_data_pointer(self, builder, ptr, align=None):
 @register_default(types.UnicodeCharSeq)
 class UnicodeCharSeq(DataModel):
     def __init__(self, dmm, fe_type):
-        super(UnicodeCharSeq, self).__init__(dmm, fe_type)
+        super().__init__(dmm, fe_type)
         charty = ir.IntType(numpy_support.sizeof_unicode_char * 8)
         self._be_type = ir.ArrayType(charty, fe_type.count)
 
@@ -1044,7 +1044,7 @@ def from_argument(self, builder, value):
 @register_default(types.CharSeq)
 class CharSeq(DataModel):
     def __init__(self, dmm, fe_type):
-        super(CharSeq, self).__init__(dmm, fe_type)
+        super().__init__(dmm, fe_type)
         charty = ir.IntType(8)
         self._be_type = ir.ArrayType(charty, fe_type.count)
 
@@ -1086,7 +1086,7 @@ def __init__(self, dmm, fe_type, need_indices):
         if need_indices:
             # For ndenumerate()
             members.append(("indices", types.EphemeralArray(types.intp, ndim)))
-        super(CContiguousFlatIter, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 class FlatIter(StructModel):
@@ -1100,7 +1100,7 @@ def __init__(self, dmm, fe_type):
             ("indices", types.EphemeralArray(types.intp, ndim)),
             ("exhausted", types.EphemeralPointer(types.boolean)),
         ]
-        super(FlatIter, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.UniTupleIter)
@@ -1113,7 +1113,7 @@ def __init__(self, dmm, fe_type):
                 fe_type.container,
             ),
         ]
-        super(UniTupleIter, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.misc.SliceLiteral)
@@ -1125,7 +1125,7 @@ def __init__(self, dmm, fe_type):
             ("stop", types.intp),
             ("step", types.intp),
         ]
-        super(SliceModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.NPDatetime)
@@ -1133,7 +1133,7 @@ def __init__(self, dmm, fe_type):
 class NPDatetimeModel(PrimitiveModel):
     def __init__(self, dmm, fe_type):
         be_type = ir.IntType(64)
-        super(NPDatetimeModel, self).__init__(dmm, fe_type, be_type)
+        super().__init__(dmm, fe_type, be_type)
 
 
 @register_default(types.ArrayIterator)
@@ -1144,7 +1144,7 @@ def __init__(self, dmm, fe_type):
             ("index", types.EphemeralPointer(types.uintp)),
             ("array", fe_type.array_type),
         ]
-        super(ArrayIterator, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.EnumerateType)
@@ -1155,7 +1155,7 @@ def __init__(self, dmm, fe_type):
             ("iter", fe_type.source_type),
         ]
 
-        super(EnumerateType, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.ZipType)
@@ -1165,7 +1165,7 @@ def __init__(self, dmm, fe_type):
             ("iter%d" % i, source_type.iterator_type)
             for i, source_type in enumerate(fe_type.source_types)
         ]
-        super(ZipType, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.RangeIteratorType)
@@ -1178,13 +1178,13 @@ def __init__(self, dmm, fe_type):
             ("step", int_type),
             ("count", types.EphemeralPointer(int_type)),
         ]
-        super(RangeIteratorType, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.Generator)
 class GeneratorModel(CompositeModel):
     def __init__(self, dmm, fe_type):
-        super(GeneratorModel, self).__init__(dmm, fe_type)
+        super().__init__(dmm, fe_type)
         # XXX Fold this in DataPacker?
         self._arg_models = [
             self._dmm.lookup(t)
@@ -1253,7 +1253,7 @@ def __init__(self, dmm, fe_type):
             ("data", types.CPointer(fe_type.dtype)),
             ("meminfo", types.MemInfoPointer(fe_type.dtype)),
         ]
-        super(ArrayCTypesModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.RangeType)
@@ -1261,7 +1261,7 @@ class RangeModel(StructModel):
     def __init__(self, dmm, fe_type):
         int_type = fe_type.iterator_type.yield_type
         members = [("start", int_type), ("stop", int_type), ("step", int_type)]
-        super(RangeModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 # =============================================================================
@@ -1276,7 +1276,7 @@ def __init__(self, dmm, fe_type):
             ("indices", types.EphemeralArray(types.intp, ndim)),
             ("exhausted", types.EphemeralPointer(types.boolean)),
         ]
-        super(NdIndexModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.NumpyFlatType)
@@ -1336,13 +1336,13 @@ def __init__(self, dmm, fe_type):
                 member_name = "scalar%d" % i
                 members.append((member_name, types.EphemeralPointer(ty)))
 
-        super(NdIter, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @register_default(types.DeferredType)
 class DeferredStructModel(CompositeModel):
     def __init__(self, dmm, fe_type):
-        super(DeferredStructModel, self).__init__(dmm, fe_type)
+        super().__init__(dmm, fe_type)
         self.typename = "deferred.{0}".format(id(fe_type))
         self.actual_fe_type = fe_type.get()
 
diff --git a/numba_cuda/numba/cuda/datamodel/cuda_packer.py b/numba_cuda/numba/cuda/datamodel/cuda_packer.py
index c879971ec..9e11d8a5d 100644
--- a/numba_cuda/numba/cuda/datamodel/cuda_packer.py
+++ b/numba_cuda/numba/cuda/datamodel/cuda_packer.py
@@ -7,7 +7,7 @@
 from numba.cuda import cgutils
 
 
-class DataPacker(object):
+class DataPacker:
     """
     A helper to pack a number of typed arguments into a data structure.
     Omitted arguments (i.e. values with the type `Omitted`) are automatically
@@ -64,7 +64,7 @@ def load_into(self, builder, ptr, formal_list):
         self._do_load(builder, ptr, formal_list)
 
 
-class ArgPacker(object):
+class ArgPacker:
     """
     Compute the position for each high-level typed argument.
     It flattens every composite argument into primitive types.
@@ -167,7 +167,7 @@ def rec(iterable):
 _POP = 4
 
 
-class _Unflattener(object):
+class _Unflattener:
     """
     An object used to unflatten nested sequences after a given pattern
     (an arbitrarily nested sequence).
diff --git a/numba_cuda/numba/cuda/datamodel/cuda_testing.py b/numba_cuda/numba/cuda/datamodel/cuda_testing.py
index 35dcc11d9..299feaea8 100644
--- a/numba_cuda/numba/cuda/datamodel/cuda_testing.py
+++ b/numba_cuda/numba/cuda/datamodel/cuda_testing.py
@@ -90,7 +90,7 @@ def test_as_return(self):
         str(materialized)
 
 
-class SupportAsDataMixin(object):
+class SupportAsDataMixin:
     """Test as_data() and from_data()"""
 
     # XXX test load_from_data_pointer() as well
@@ -119,7 +119,7 @@ def test_as_data(self):
         str(materialized)
 
 
-class NotSupportAsDataMixin(object):
+class NotSupportAsDataMixin:
     """Ensure as_data() and from_data() raise NotImplementedError."""
 
     def test_as_data_not_supported(self):
diff --git a/numba_cuda/numba/cuda/deviceufunc.py b/numba_cuda/numba/cuda/deviceufunc.py
index ca6be4bd4..f69093ee0 100644
--- a/numba_cuda/numba/cuda/deviceufunc.py
+++ b/numba_cuda/numba/cuda/deviceufunc.py
@@ -136,7 +136,7 @@ def _multi_broadcast(*shapelist):
         return result
 
 
-class UFuncMechanism(object):
+class UFuncMechanism:
     """
     Prepare ufunc arguments for vectorize.
     """
@@ -635,7 +635,7 @@ def _gen_src_index(adims, atype):
         return "__tid__"
 
 
-class GUFuncEngine(object):
+class GUFuncEngine:
     """Determine how to broadcast and execute a gufunc
     base on input shape and signature
     """
@@ -707,7 +707,7 @@ def schedule(self, ishapes):
         return GUFuncSchedule(self, inner_shapes, oshapes, loopdims, pinned)
 
 
-class GUFuncSchedule(object):
+class GUFuncSchedule:
     def __init__(self, parent, ishapes, oshapes, loopdims, pinned):
         self.parent = parent
         # core shapes
@@ -729,7 +729,7 @@ def __str__(self):
         return pprint.pformat(dict(values))
 
 
-class GeneralizedUFunc(object):
+class GeneralizedUFunc:
     def __init__(self, kernelmap, engine):
         self.kernelmap = kernelmap
         self.engine = engine
diff --git a/numba_cuda/numba/cuda/dispatcher.py b/numba_cuda/numba/cuda/dispatcher.py
index c0ee303a4..aa8fb76be 100644
--- a/numba_cuda/numba/cuda/dispatcher.py
+++ b/numba_cuda/numba/cuda/dispatcher.py
@@ -638,7 +638,7 @@ def _prepare_args(self, ty, val, stream, retr, kernelargs):
             raise NotImplementedError(ty, val)
 
 
-class ForAll(object):
+class ForAll:
     def __init__(self, dispatcher, ntasks, tpb, stream, sharedmem):
         if ntasks < 0:
             raise ValueError(
@@ -762,7 +762,7 @@ def load_overload(self, sig, target_context):
             return super().load_overload(sig, target_context)
 
 
-class OmittedArg(object):
+class OmittedArg:
     """
     A placeholder for omitted arguments with a default value.
     """
@@ -778,7 +778,7 @@ def _numba_type_(self):
         return types.Omitted(self.value)
 
 
-class CompilingCounter(object):
+class CompilingCounter:
     """
     A simple counter that increment in __enter__ and decrement in __exit__.
     """
@@ -1391,7 +1391,7 @@ def _set_uuid(self, u):
 )
 
 
-class _FunctionCompiler(object):
+class _FunctionCompiler:
     def __init__(self, py_func, targetdescr, targetoptions, pipeline_class):
         self.py_func = py_func
         self.targetdescr = targetdescr
@@ -2414,7 +2414,7 @@ def cb_llvm(dur):
 class ObjModeLiftedWith(LiftedWith):
     def __init__(self, *args, **kwargs):
         self.output_types = kwargs.pop("output_types", None)
-        super(LiftedWith, self).__init__(*args, **kwargs)
+        super().__init__(*args, **kwargs)
         if not self.flags.force_pyobject:
             raise ValueError("expecting `flags.force_pyobject`")
         if self.output_types is None:
diff --git a/numba_cuda/numba/cuda/errors.py b/numba_cuda/numba/cuda/errors.py
index 4320f3b30..5f6f5c8ee 100644
--- a/numba_cuda/numba/cuda/errors.py
+++ b/numba_cuda/numba/cuda/errors.py
@@ -12,7 +12,7 @@ def __init__(self, msg, tid=None, ctaid=None):
         self.msg = msg
         t = "An exception was raised in thread=%s block=%s\n\t%s"
         msg = t % (self.tid, self.ctaid, self.msg)
-        super(KernelRuntimeError, self).__init__(msg)
+        super().__init__(msg)
 
 
 class UnsupportedBytecodeError(Exception):
diff --git a/numba_cuda/numba/cuda/kernels/reduction.py b/numba_cuda/numba/cuda/kernels/reduction.py
index 463db8846..da0a555eb 100644
--- a/numba_cuda/numba/cuda/kernels/reduction.py
+++ b/numba_cuda/numba/cuda/kernels/reduction.py
@@ -166,7 +166,7 @@ def gpu_reduce_block_strided(arr, partials, init, use_init):
     return cuda.jit(gpu_reduce_block_strided)
 
 
-class Reduce(object):
+class Reduce:
     """Create a reduction object that reduces values using a given binary
     function. The binary function is compiled once and cached inside this
     object. Keeping this object alive will prevent re-compilation.
diff --git a/numba_cuda/numba/cuda/lowering.py b/numba_cuda/numba/cuda/lowering.py
index f97447cc4..c0b18b40a 100644
--- a/numba_cuda/numba/cuda/lowering.py
+++ b/numba_cuda/numba/cuda/lowering.py
@@ -38,7 +38,7 @@
 _VarArgItem = namedtuple("_VarArgItem", ("vararg", "index"))
 
 
-class BaseLower(object):
+class BaseLower:
     """
     Lower IR to LLVM
     """
@@ -467,7 +467,7 @@ def _find_singly_assigned_variable(self):
     def pre_block(self, block):
         from numba.cuda.core.unsafe import eh
 
-        super(Lower, self).pre_block(block)
+        super().pre_block(block)
         self._cur_ir_block = block
 
         if block == self.firstblk:
diff --git a/numba_cuda/numba/cuda/memory_management/nrt_context.py b/numba_cuda/numba/cuda/memory_management/nrt_context.py
index 208412d3b..3e171779a 100644
--- a/numba_cuda/numba/cuda/memory_management/nrt_context.py
+++ b/numba_cuda/numba/cuda/memory_management/nrt_context.py
@@ -32,7 +32,7 @@
 meminfo_data_ty = ir.FunctionType(_pointer_type, [_pointer_type])
 
 
-class NRTContext(object):
+class NRTContext:
     """
     An object providing access to NRT APIs in the lowering pass.
     """
diff --git a/numba_cuda/numba/cuda/misc/appdirs.py b/numba_cuda/numba/cuda/misc/appdirs.py
index 78a660702..09c017561 100644
--- a/numba_cuda/numba/cuda/misc/appdirs.py
+++ b/numba_cuda/numba/cuda/misc/appdirs.py
@@ -364,7 +364,7 @@ def user_log_dir(appname=None, appauthor=None, version=None, opinion=True):
     return path
 
 
-class AppDirs(object):
+class AppDirs:
     """Convenience wrapper for getting application dirs."""
 
     def __init__(
diff --git a/numba_cuda/numba/cuda/misc/special.py b/numba_cuda/numba/cuda/misc/special.py
index 2bf1a2342..2171721ed 100644
--- a/numba_cuda/numba/cuda/misc/special.py
+++ b/numba_cuda/numba/cuda/misc/special.py
@@ -2,7 +2,7 @@
 # SPDX-License-Identifier: BSD-2-Clause
 
 
-class prange(object):
+class prange:
     """Provides a 1D parallel iterator that generates a sequence of integers.
     In non-parallel contexts, prange is identical to range.
     """
diff --git a/numba_cuda/numba/cuda/models.py b/numba_cuda/numba/cuda/models.py
index 88a91e9e1..862ff3a8f 100644
--- a/numba_cuda/numba/cuda/models.py
+++ b/numba_cuda/numba/cuda/models.py
@@ -43,7 +43,7 @@ def __init__(self, dmm, fe_type):
             be_type = ir.DoubleType()
         else:
             raise NotImplementedError(fe_type)
-        super(FloatModel, self).__init__(dmm, fe_type, be_type)
+        super().__init__(dmm, fe_type, be_type)
 
 
 register_model(CUDADispatcher)(models.OpaqueModel)
@@ -53,4 +53,4 @@ def __init__(self, dmm, fe_type):
 class _model___nv_bfloat16(PrimitiveModel):
     def __init__(self, dmm, fe_type):
         be_type = ir.IntType(16)
-        super(_model___nv_bfloat16, self).__init__(dmm, fe_type, be_type)
+        super().__init__(dmm, fe_type, be_type)
diff --git a/numba_cuda/numba/cuda/np/arrayobj.py b/numba_cuda/numba/cuda/np/arrayobj.py
index 2922a73f0..115fb63f2 100644
--- a/numba_cuda/numba/cuda/np/arrayobj.py
+++ b/numba_cuda/numba/cuda/np/arrayobj.py
@@ -118,7 +118,7 @@ def _make_refs(self, ref):
             try:
                 array_impl = self._context.get_function("__array__", sig)
             except NotImplementedError:
-                return super(ArrayStruct, self)._make_refs(ref)
+                return super()._make_refs(ref)
 
             # Return a wrapped structure and its unwrapped reference
             datamodel = self._context.data_model_manager[array_type]
@@ -701,7 +701,7 @@ def array_itemset(context, builder, sig, args):
 # Advanced / fancy indexing
 
 
-class Indexer(object):
+class Indexer:
     """
     Generic indexer interface, for generating indices over a fancy indexed
     array on a single dimension.
@@ -1061,7 +1061,7 @@ def loop_tail(self):
         builder.position_at_end(self.bb_end)
 
 
-class FancyIndexer(object):
+class FancyIndexer:
     """
     Perform fancy indexing on the given array.
     """
@@ -3827,7 +3827,7 @@ def make_nditer_cls(nditerty):
     narrays = len(nditerty.arrays)
     nshapes = ndim if nditerty.need_shaped_indexing else 1
 
-    class BaseSubIter(object):
+    class BaseSubIter:
         """
         Base class for sub-iterators of a nditer() instance.
         """
diff --git a/numba_cuda/numba/cuda/np/npyimpl.py b/numba_cuda/numba/cuda/np/npyimpl.py
index 5b1a5d275..a25779d9a 100644
--- a/numba_cuda/numba/cuda/np/npyimpl.py
+++ b/numba_cuda/numba/cuda/np/npyimpl.py
@@ -52,7 +52,7 @@
 # generated and reading performs the appropriate indirection.
 
 
-class _ScalarIndexingHelper(object):
+class _ScalarIndexingHelper:
     def update_indices(self, loop_indices, name):
         pass
 
@@ -60,7 +60,7 @@ def as_values(self):
         pass
 
 
-class _ScalarHelper(object):
+class _ScalarHelper:
     """Helper class to handle scalar arguments (and result).
     Note that store_data is only used when generating code for
     a scalar ufunc and to write the output value.
@@ -764,7 +764,7 @@ def numpy_gufunc_kernel(context, builder, sig, args, ufunc, kernel_class):
 
 
 # Kernels are the code to be executed inside the multidimensional loop.
-class _Kernel(object):
+class _Kernel:
     def __init__(self, context, builder, outer_sig):
         self.context = context
         self.builder = builder
@@ -842,7 +842,7 @@ def _ufunc_db_function(ufunc):
 
     class _KernelImpl(_Kernel):
         def __init__(self, context, builder, outer_sig):
-            super(_KernelImpl, self).__init__(context, builder, outer_sig)
+            super().__init__(context, builder, outer_sig)
             loop = ufunc_find_matching_loop(
                 ufunc,
                 outer_sig.args + tuple(_unpack_output_types(ufunc, outer_sig)),
diff --git a/numba_cuda/numba/cuda/np/polynomial/polynomial_core.py b/numba_cuda/numba/cuda/np/polynomial/polynomial_core.py
index a8a28f8f2..d7b497abe 100644
--- a/numba_cuda/numba/cuda/np/polynomial/polynomial_core.py
+++ b/numba_cuda/numba/cuda/np/polynomial/polynomial_core.py
@@ -37,7 +37,7 @@ def __init__(self, dmm, fe_type):
             # Introduced in NumPy 1.24, maybe leave it out for now
             # ('symbol', types.string)
         ]
-        super(PolynomialModel, self).__init__(dmm, fe_type, members)
+        super().__init__(dmm, fe_type, members)
 
 
 @type_callable(Polynomial)
diff --git a/numba_cuda/numba/cuda/np/ufunc/decorators.py b/numba_cuda/numba/cuda/np/ufunc/decorators.py
index f71cecbf5..d9e0e1966 100644
--- a/numba_cuda/numba/cuda/np/ufunc/decorators.py
+++ b/numba_cuda/numba/cuda/np/ufunc/decorators.py
@@ -5,7 +5,7 @@
 from numba.cuda.vectorizers import CUDAVectorize, CUDAGUFuncVectorize
 
 
-class _BaseVectorize(object):
+class _BaseVectorize:
     @classmethod
     def get_identity(cls, kwargs):
         return kwargs.pop("identity", None)
diff --git a/numba_cuda/numba/cuda/np/ufunc/ufuncbuilder.py b/numba_cuda/numba/cuda/np/ufunc/ufuncbuilder.py
index c12e38c52..257ab37d6 100644
--- a/numba_cuda/numba/cuda/np/ufunc/ufuncbuilder.py
+++ b/numba_cuda/numba/cuda/np/ufunc/ufuncbuilder.py
@@ -24,7 +24,7 @@ def _compile_element_wise_function(nb_func, targetoptions, sig):
 # Class definitions
 
 
-class _BaseUFuncBuilder(object):
+class _BaseUFuncBuilder:
     def add(self, sig=None):
         if hasattr(self, "targetoptions"):
             targetoptions = self.targetoptions
diff --git a/numba_cuda/numba/cuda/nvvmutils.py b/numba_cuda/numba/cuda/nvvmutils.py
index 56a7dd01a..f6c4a2b4e 100644
--- a/numba_cuda/numba/cuda/nvvmutils.py
+++ b/numba_cuda/numba/cuda/nvvmutils.py
@@ -219,7 +219,7 @@ def call_sreg(builder, name):
     return builder.call(fn, ())
 
 
-class SRegBuilder(object):
+class SRegBuilder:
     def __init__(self, builder):
         self.builder = builder
 
diff --git a/numba_cuda/numba/cuda/simulator/api.py b/numba_cuda/numba/cuda/simulator/api.py
index 5851471bd..de070f310 100644
--- a/numba_cuda/numba/cuda/simulator/api.py
+++ b/numba_cuda/numba/cuda/simulator/api.py
@@ -37,7 +37,7 @@ def is_bfloat16_supported():
     return False
 
 
-class stream(object):
+class stream:
     """
     The stream API is supported in the simulator - however, all execution
     occurs synchronously, so synchronization requires no operation.
@@ -94,7 +94,7 @@ def get_current_device():
 # Events
 
 
-class Event(object):
+class Event:
     """
     The simulator supports the event API, but they do not record timing info,
     and all simulation is synchronous. Execution time is not recorded.
diff --git a/numba_cuda/numba/cuda/simulator/compiler.py b/numba_cuda/numba/cuda/simulator/compiler.py
index d2fb21fea..11f5f31d8 100644
--- a/numba_cuda/numba/cuda/simulator/compiler.py
+++ b/numba_cuda/numba/cuda/simulator/compiler.py
@@ -18,7 +18,7 @@ def run_frontend(func):
     pass
 
 
-class DefaultPassBuilder(object):
+class DefaultPassBuilder:
     @staticmethod
     def define_nopython_lowering_pipeline(state, name="nopython_lowering"):
         pass
diff --git a/numba_cuda/numba/cuda/simulator/cudadrv/devicearray.py b/numba_cuda/numba/cuda/simulator/cudadrv/devicearray.py
index 60103ea43..76c3696a1 100644
--- a/numba_cuda/numba/cuda/simulator/cudadrv/devicearray.py
+++ b/numba_cuda/numba/cuda/simulator/cudadrv/devicearray.py
@@ -36,10 +36,10 @@ class FakeShape(tuple):
     def __getitem__(self, k):
         if isinstance(k, int) and k < 0:
             raise IndexError("tuple index out of range")
-        return super(FakeShape, self).__getitem__(k)
+        return super().__getitem__(k)
 
 
-class FakeWithinKernelCUDAArray(object):
+class FakeWithinKernelCUDAArray:
     """
     Created to emulate the behavior of arrays within kernels, where either
     array.item or array['item'] is valid (that is, give all structured
@@ -99,7 +99,7 @@ def convert_fakes(obj):
         return call(*args, **kwargs)
 
 
-class FakeCUDAArray(object):
+class FakeCUDAArray:
     """
     Implements the interface of a DeviceArray/DeviceRecord, but mostly just
     wraps a NumPy array.
diff --git a/numba_cuda/numba/cuda/simulator/cudadrv/driver.py b/numba_cuda/numba/cuda/simulator/cudadrv/driver.py
index 679112cab..25d750a3c 100644
--- a/numba_cuda/numba/cuda/simulator/cudadrv/driver.py
+++ b/numba_cuda/numba/cuda/simulator/cudadrv/driver.py
@@ -27,7 +27,7 @@ def device_to_device(dst, src, size, stream=0):
     host_to_device(dst, src, size)
 
 
-class FakeDriver(object):
+class FakeDriver:
     def get_device_count(self):
         return 1
 
diff --git a/numba_cuda/numba/cuda/simulator/cudadrv/nvvm.py b/numba_cuda/numba/cuda/simulator/cudadrv/nvvm.py
index b114eba9c..1d067ae30 100644
--- a/numba_cuda/numba/cuda/simulator/cudadrv/nvvm.py
+++ b/numba_cuda/numba/cuda/simulator/cudadrv/nvvm.py
@@ -11,7 +11,7 @@ class NvvmSupportError(ImportError):
     pass
 
 
-class NVVM(object):
+class NVVM:
     def __init__(self):
         raise NvvmSupportError("NVVM not supported in the simulator")
 
diff --git a/numba_cuda/numba/cuda/simulator/cudadrv/runtime.py b/numba_cuda/numba/cuda/simulator/cudadrv/runtime.py
index 353a9e0a7..50ebe8838 100644
--- a/numba_cuda/numba/cuda/simulator/cudadrv/runtime.py
+++ b/numba_cuda/numba/cuda/simulator/cudadrv/runtime.py
@@ -7,7 +7,7 @@
 """
 
 
-class FakeRuntime(object):
+class FakeRuntime:
     def get_version(self):
         return (-1, -1)
 
diff --git a/numba_cuda/numba/cuda/simulator/kernel.py b/numba_cuda/numba/cuda/simulator/kernel.py
index 4c60971d6..7c29ecf9a 100644
--- a/numba_cuda/numba/cuda/simulator/kernel.py
+++ b/numba_cuda/numba/cuda/simulator/kernel.py
@@ -61,7 +61,7 @@ def __getitem__(self, key):
         return FakeOverload()
 
 
-class FakeCUDAKernel(object):
+class FakeCUDAKernel:
     """
     Wraps a @cuda.jit-ed function.
     """
@@ -181,7 +181,7 @@ def debug_wrapper(*args, **kwargs):
         else:
             target = f
 
-        super(BlockThread, self).__init__(target=target)
+        super().__init__(target=target)
         self.syncthreads_event = threading.Event()
         self.syncthreads_blocked = False
         self._manager = manager
@@ -198,7 +198,7 @@ def debug_wrapper(*args, **kwargs):
 
     def run(self):
         try:
-            super(BlockThread, self).run()
+            super().run()
         except Exception as e:
             tid = "tid=%s" % list(self.threadIdx)
             ctaid = "ctaid=%s" % list(self.blockIdx)
@@ -250,7 +250,7 @@ def __str__(self):
         return "Thread <<<%s, %s>>>" % (self.blockIdx, self.threadIdx)
 
 
-class BlockManager(object):
+class BlockManager:
     """
     Manages the execution of a thread block.
 
diff --git a/numba_cuda/numba/cuda/simulator/kernelapi.py b/numba_cuda/numba/cuda/simulator/kernelapi.py
index bed1cdc20..aed3bed85 100644
--- a/numba_cuda/numba/cuda/simulator/kernelapi.py
+++ b/numba_cuda/numba/cuda/simulator/kernelapi.py
@@ -19,7 +19,7 @@
 from .vector_types import vector_types
 
 
-class Dim3(object):
+class Dim3:
     """
     Used to implement thread/block indices/dimensions
     """
@@ -62,7 +62,7 @@ def this_grid(self):
         return GridGroup()
 
 
-class FakeCUDALocal(object):
+class FakeCUDALocal:
     """
     CUDA Local arrays
     """
@@ -76,7 +76,7 @@ def array(self, shape, dtype, alignment=None):
         return np.empty(shape, dtype)
 
 
-class FakeCUDAConst(object):
+class FakeCUDAConst:
     """
     CUDA Const arrays
     """
@@ -85,7 +85,7 @@ def array_like(self, ary):
         return ary
 
 
-class FakeCUDAShared(object):
+class FakeCUDAShared:
     """
     CUDA Shared arrays.
 
@@ -150,7 +150,7 @@ def array(self, shape, dtype, alignment=None):
 exchlock = threading.Lock()
 
 
-class FakeCUDAAtomic(object):
+class FakeCUDAAtomic:
     def add(self, array, index, val):
         with addlock:
             old = array[index]
@@ -245,7 +245,7 @@ def cas(self, array, index, old, val):
             return loaded
 
 
-class FakeCUDAFp16(object):
+class FakeCUDAFp16:
     def hadd(self, a, b):
         return a + b
 
@@ -337,7 +337,7 @@ def hmin(self, a, b):
         return min(a, b)
 
 
-class FakeCUDAModule(object):
+class FakeCUDAModule:
     """
     An instance of this class will be injected into the __globals__ for an
     executing function in order to implement calls to cuda.*. This will fail to
diff --git a/numba_cuda/numba/cuda/stubs.py b/numba_cuda/numba/cuda/stubs.py
index b15a84097..82c0fb24b 100644
--- a/numba_cuda/numba/cuda/stubs.py
+++ b/numba_cuda/numba/cuda/stubs.py
@@ -12,7 +12,7 @@
 from inspect import Signature, Parameter
 
 
-class Stub(object):
+class Stub:
     """
     A stub object to represent special objects that are meaningless
     outside the context of a CUDA kernel
diff --git a/numba_cuda/numba/cuda/target.py b/numba_cuda/numba/cuda/target.py
index 1f2ac98c9..8e2c9a4fb 100644
--- a/numba_cuda/numba/cuda/target.py
+++ b/numba_cuda/numba/cuda/target.py
@@ -89,7 +89,7 @@ def resolve_value_type(self, val):
                     val = disp
 
         # continue with parent logic
-        return super(CUDATypingContext, self).resolve_value_type(val)
+        return super().resolve_value_type(val)
 
     def can_convert(self, fromty, toty):
         """
diff --git a/numba_cuda/numba/cuda/testing.py b/numba_cuda/numba/cuda/testing.py
index 7906af27b..d9a295059 100644
--- a/numba_cuda/numba/cuda/testing.py
+++ b/numba_cuda/numba/cuda/testing.py
@@ -370,7 +370,7 @@ def skip_on_nvjitlink_13_1_sm_120(reason):
     return unittest.skipIf(_is_nvjitlink_13_1_and_sm_120(), reason)
 
 
-class ForeignArray(object):
+class ForeignArray:
     """
     Class for emulating an array coming from another library through the CUDA
     Array interface. This just hides a DeviceNDArray so that it doesn't look
diff --git a/numba_cuda/numba/cuda/tests/core/test_serialize.py b/numba_cuda/numba/cuda/tests/core/test_serialize.py
index d2da82c47..a2798b4a0 100644
--- a/numba_cuda/numba/cuda/tests/core/test_serialize.py
+++ b/numba_cuda/numba/cuda/tests/core/test_serialize.py
@@ -310,7 +310,7 @@ def test_dynamic_class_issue_7356(self):
         self.assertEqual(cfunc(), (100, 100))
 
 
-class DynClass(object):
+class DynClass:
     # For testing issue #7356
     a = None
 
diff --git a/numba_cuda/numba/cuda/tests/cudadrv/test_cuda_memory.py b/numba_cuda/numba/cuda/tests/cudadrv/test_cuda_memory.py
index d6852c8c5..4ccd6b6be 100644
--- a/numba_cuda/numba/cuda/tests/cudadrv/test_cuda_memory.py
+++ b/numba_cuda/numba/cuda/tests/cudadrv/test_cuda_memory.py
@@ -19,7 +19,7 @@ def setUp(self):
     def tearDown(self):
         self.context.reset()
         del self.context
-        super(TestCudaMemory, self).tearDown()
+        super().tearDown()
 
     def _template(self, obj):
         self.assertTrue(driver.is_device_memory(obj))
@@ -112,7 +112,7 @@ def setUp(self):
 
     def tearDown(self):
         del self.context
-        super(TestCudaMemoryFunctions, self).tearDown()
+        super().tearDown()
 
     def test_memcpy(self):
         hstary = np.arange(100, dtype=np.uint32)
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_array_methods.py b/numba_cuda/numba/cuda/tests/cudapy/test_array_methods.py
index f30d42de7..4921242a4 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_array_methods.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_array_methods.py
@@ -18,11 +18,11 @@ class TestCudaArrayMethods(CUDATestCase):
     def setUp(self):
         self.old_nrt_setting = config.CUDA_ENABLE_NRT
         config.CUDA_ENABLE_NRT = True
-        super(TestCudaArrayMethods, self).setUp()
+        super().setUp()
 
     def tearDown(self):
         config.CUDA_ENABLE_NRT = self.old_nrt_setting
-        super(TestCudaArrayMethods, self).tearDown()
+        super().tearDown()
 
     def test_reinterpret_array_type(self):
         """
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_array_reductions.py b/numba_cuda/numba/cuda/tests/cudapy/test_array_reductions.py
index dddcbfd4b..0ead25770 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_array_reductions.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_array_reductions.py
@@ -16,7 +16,7 @@ class TestArrayReductions(MemoryLeakMixin, TestCase):
     """
 
     def setUp(self):
-        super(TestArrayReductions, self).setUp()
+        super().setUp()
         np.random.seed(42)
         self.old_nrt_setting = config.CUDA_ENABLE_NRT
         self.old_perf_warnings_setting = config.DISABLE_PERFORMANCE_WARNINGS
@@ -26,7 +26,7 @@ def setUp(self):
     def tearDown(self):
         config.CUDA_ENABLE_NRT = self.old_nrt_setting
         config.DISABLE_PERFORMANCE_WARNINGS = self.old_perf_warnings_setting
-        super(TestArrayReductions, self).tearDown()
+        super().tearDown()
 
     def test_all_basic(self):
         cases = (
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_extending_types.py b/numba_cuda/numba/cuda/tests/cudapy/test_extending_types.py
index 60b465997..dfcd8487e 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_extending_types.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_extending_types.py
@@ -37,7 +37,7 @@ class TestExtTypDummy(unittest.TestCase):
     def setUp(self):
         class DummyType(types.Type):
             def __init__(self):
-                super(DummyType, self).__init__(name="Dummy")
+                super().__init__(name="Dummy")
 
         make_attribute_wrapper(DummyType, "value", "value")
 
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_retrieve_autoconverted_arrays.py b/numba_cuda/numba/cuda/tests/cudapy/test_retrieve_autoconverted_arrays.py
index 30db2844c..fd2bd3590 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_retrieve_autoconverted_arrays.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_retrieve_autoconverted_arrays.py
@@ -9,7 +9,7 @@
 import unittest
 
 
-class DefaultIn(object):
+class DefaultIn:
     def prepare_args(self, ty, val, **kwargs):
         return ty, wrap_arg(val, default=cuda.In)
 
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_tracing.py b/numba_cuda/numba/cuda/tests/cudapy/test_tracing.py
index d2234f150..fac093d51 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_tracing.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_tracing.py
@@ -39,7 +39,7 @@ def getvalue(self):
         return log
 
 
-class Class(object):
+class Class:
     @tracing.trace
     @classmethod
     def class_method(cls):
@@ -69,7 +69,7 @@ def __repr__(self):
         return "<Class instance>"
 
 
-class Class2(object):
+class Class2:
     @classmethod
     def class_method(cls):
         pass
@@ -118,7 +118,7 @@ def test_traced_function():
 
 class TestTracing(unittest.TestCase):
     def __init__(self, *args):
-        super(TestTracing, self).__init__(*args)
+        super().__init__(*args)
 
     def setUp(self):
         self.capture = CapturedTrace()
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_vectorize_decor.py b/numba_cuda/numba/cuda/tests/cudapy/test_vectorize_decor.py
index d9d7c78ec..b2b8b2015 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_vectorize_decor.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_vectorize_decor.py
@@ -33,7 +33,7 @@ def vector_add(a, b):
     return a + b
 
 
-class BaseVectorizeDecor(object):
+class BaseVectorizeDecor:
     target = None
     wrapper = None
     funcs = {
diff --git a/numba_cuda/numba/cuda/tests/nrt/test_nrt_refct.py b/numba_cuda/numba/cuda/tests/nrt/test_nrt_refct.py
index 8256f8a40..11af80f0f 100644
--- a/numba_cuda/numba/cuda/tests/nrt/test_nrt_refct.py
+++ b/numba_cuda/numba/cuda/tests/nrt/test_nrt_refct.py
@@ -14,17 +14,17 @@
 @skip_on_cudasim("No refcounting in the simulator")
 class TestNrtRefCt(EnableNRTStatsMixin, CUDATestCase):
     def setUp(self):
-        super(TestNrtRefCt, self).setUp()
+        super().setUp()
 
     def tearDown(self):
-        super(TestNrtRefCt, self).tearDown()
+        super().tearDown()
 
     def run(self, result=None):
         with (
             override_config("CUDA_ENABLE_NRT", True),
             override_config("CUDA_NRT_STATS", True),
         ):
-            super(TestNrtRefCt, self).run(result)
+            super().run(result)
 
     def test_no_return(self):
         """
diff --git a/numba_cuda/numba/cuda/tests/support.py b/numba_cuda/numba/cuda/tests/support.py
index 8d7a294b8..24be3c07b 100644
--- a/numba_cuda/numba/cuda/tests/support.py
+++ b/numba_cuda/numba/cuda/tests/support.py
@@ -47,7 +47,7 @@
     from numba.core import types as upstream_types
 
 
-class EnableNRTStatsMixin(object):
+class EnableNRTStatsMixin:
     """Mixin to enable the NRT statistics counters."""
 
     def setUp(self):
@@ -762,7 +762,7 @@ def make_dummy_type(self):
         dummy_type = DummyType("my_dummy")
         register_model(DummyType)(OpaqueModel)
 
-        class Dummy(object):
+        class Dummy:
             pass
 
         @typeof_impl.register(Dummy)
@@ -787,7 +787,7 @@ def unbox_dummy(typ, obj, c):
         return Dummy, DummyType
 
 
-class MemoryLeak(object):
+class MemoryLeak:
     __enable_leak_check = True
 
     def memory_leak_setup(self):
@@ -816,16 +816,16 @@ def disable_leak_check(self):
 
 class MemoryLeakMixin(EnableNRTStatsMixin, MemoryLeak):
     def setUp(self):
-        super(MemoryLeakMixin, self).setUp()
+        super().setUp()
         self.memory_leak_setup()
 
     def tearDown(self):
         gc.collect()
         self.memory_leak_teardown()
-        super(MemoryLeakMixin, self).tearDown()
+        super().tearDown()
 
 
-class CheckWarningsMixin(object):
+class CheckWarningsMixin:
     @contextlib.contextmanager
     def check_warnings(self, messages, category=RuntimeWarning):
         with warnings.catch_warnings(record=True) as catch:
diff --git a/numba_cuda/numba/cuda/typeconv/castgraph.py b/numba_cuda/numba/cuda/typeconv/castgraph.py
index aa9f20db7..7dbe37f81 100644
--- a/numba_cuda/numba/cuda/typeconv/castgraph.py
+++ b/numba_cuda/numba/cuda/typeconv/castgraph.py
@@ -30,7 +30,7 @@ class Conversion(enum.IntEnum):
     nil = 99
 
 
-class CastSet(object):
+class CastSet:
     """A set of casting rules.
 
     There is at most one rule per target type.
@@ -71,7 +71,7 @@ def __getitem__(self, item):
         return self._rels[item]
 
 
-class TypeGraph(object):
+class TypeGraph:
     """A graph that maintains the casting relationship of all types.
 
     This simplifies the definition of casting rules by automatically
diff --git a/numba_cuda/numba/cuda/typeconv/typeconv.py b/numba_cuda/numba/cuda/typeconv/typeconv.py
index 7510700f2..cb3cbd8f0 100644
--- a/numba_cuda/numba/cuda/typeconv/typeconv.py
+++ b/numba_cuda/numba/cuda/typeconv/typeconv.py
@@ -6,7 +6,7 @@
 from numba.cuda import types
 
 
-class TypeManager(object):
+class TypeManager:
     # The character codes used by the C/C++ API (_typeconv.cpp)
     _conversion_codes = {
         Conversion.safe: ord("s"),
@@ -59,7 +59,7 @@ def get_pointer(self):
         return _typeconv.get_pointer(self._ptr)
 
 
-class TypeCastingRules(object):
+class TypeCastingRules:
     """
     A helper for establishing type casting rules.
     """
diff --git a/numba_cuda/numba/cuda/types/cuda_abstract.py b/numba_cuda/numba/cuda/types/cuda_abstract.py
index 470f9966d..7a58d887a 100644
--- a/numba_cuda/numba/cuda/types/cuda_abstract.py
+++ b/numba_cuda/numba/cuda/types/cuda_abstract.py
@@ -134,7 +134,7 @@ def __ne__(self, other):
         return not (self == other)
 
     def __reduce__(self):
-        reconstructor, args, state = super(Type, self).__reduce__()
+        reconstructor, args, state = super().__reduce__()
         return (_type_reconstructor, (reconstructor, args, state))
 
     def unify(self, typingctx, other):
@@ -353,7 +353,7 @@ class IteratorType(IterableType):
     """
 
     def __init__(self, name, **kwargs):
-        super(IteratorType, self).__init__(name, **kwargs)
+        super().__init__(name, **kwargs)
 
     @abstractproperty
     def yield_type(self):
@@ -448,7 +448,7 @@ def __init__(self, value):
             )
         self._literal_init(value)
         fmt = "Literal[{}]({})"
-        super(Literal, self).__init__(fmt.format(type(value).__name__, value))
+        super().__init__(fmt.format(type(value).__name__, value))
 
     def _literal_init(self, value):
         self._literal_value = value
@@ -493,14 +493,14 @@ class TypeRef(Dummy):
 
     def __init__(self, instance_type):
         self.instance_type = instance_type
-        super(TypeRef, self).__init__("typeref[{}]".format(self.instance_type))
+        super().__init__("typeref[{}]".format(self.instance_type))
 
     @property
     def key(self):
         return self.instance_type
 
 
-class InitialValue(object):
+class InitialValue:
     """
     Used as a mixin for a type will potentially have an initial value that will
     be carried in the .initial_value attribute.
@@ -524,7 +524,7 @@ class Poison(Type):
 
     def __init__(self, ty):
         self.ty = ty
-        super(Poison, self).__init__(name="Poison<%s>" % ty)
+        super().__init__(name="Poison<%s>" % ty)
 
     def __unliteral__(self):
         return Poison(self)
diff --git a/numba_cuda/numba/cuda/types/cuda_common.py b/numba_cuda/numba/cuda/types/cuda_common.py
index 8e2b4c4b6..671531eeb 100644
--- a/numba_cuda/numba/cuda/types/cuda_common.py
+++ b/numba_cuda/numba/cuda/types/cuda_common.py
@@ -18,7 +18,7 @@ class Opaque(Dummy):
 class SimpleIterableType(IterableType):
     def __init__(self, name, iterator_type):
         self._iterator_type = iterator_type
-        super(SimpleIterableType, self).__init__(name)
+        super().__init__(name)
 
     @property
     def iterator_type(self):
@@ -28,7 +28,7 @@ def iterator_type(self):
 class SimpleIteratorType(IteratorType):
     def __init__(self, name, yield_type):
         self._yield_type = yield_type
-        super(SimpleIteratorType, self).__init__(name)
+        super().__init__(name)
 
     @property
     def yield_type(self):
@@ -70,7 +70,7 @@ def __init__(self, dtype, ndim, layout, readonly=False, name=None):
             if readonly:
                 type_name = "readonly %s" % type_name
             name = "%s(%s, %sd, %s)" % (type_name, dtype, ndim, layout)
-        super(Buffer, self).__init__(name)
+        super().__init__(name)
 
     @property
     def iterator_type(self):
diff --git a/numba_cuda/numba/cuda/types/cuda_containers.py b/numba_cuda/numba/cuda/types/cuda_containers.py
index edb24aceb..37884b46c 100644
--- a/numba_cuda/numba/cuda/types/cuda_containers.py
+++ b/numba_cuda/numba/cuda/types/cuda_containers.py
@@ -38,7 +38,7 @@ def __init__(self, first_type, second_type):
         self.first_type = first_type
         self.second_type = second_type
         name = "pair<%s, %s>" % (first_type, second_type)
-        super(Pair, self).__init__(name=name)
+        super().__init__(name=name)
 
     @property
     def key(self):
@@ -64,7 +64,7 @@ def __init__(self, container):
         self.container = container
         yield_type = container.dtype
         name = "iter(%s)" % container
-        super(BaseContainerIterator, self).__init__(name, yield_type)
+        super().__init__(name, yield_type)
 
     def unify(self, typingctx, other):
         cls = type(self)
@@ -89,7 +89,7 @@ def __init__(self, container):
         assert isinstance(container, self.container_class)
         self.container = container
         name = "payload(%s)" % container
-        super(BaseContainerPayload, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
@@ -240,7 +240,7 @@ def __init__(self, dtype, count):
             dtype,
             count,
         )
-        super(UniTuple, self).__init__(name)
+        super().__init__(name)
 
     @property
     def mangling_args(self):
@@ -299,7 +299,7 @@ class UnionType(Type):
     def __init__(self, types):
         self.types = tuple(sorted(set(types), key=lambda x: x.name))
         name = "Union[{}]".format(",".join(map(str, self.types)))
-        super(UnionType, self).__init__(name=name)
+        super().__init__(name=name)
 
     def get_type_tag(self, typ):
         return self.types.index(typ)
@@ -326,7 +326,7 @@ def __init__(self, types):
             self.__class__.__name__,
             ", ".join(str(i) for i in self.types),
         )
-        super(Tuple, self).__init__(name)
+        super().__init__(name)
 
     @property
     def mangling_args(self):
@@ -390,7 +390,7 @@ def __init__(self, dtype, count, cls):
         self.fields = tuple(cls._fields)
         self.instance_class = cls
         name = "%s(%s x %d)" % (cls.__name__, dtype, count)
-        super(NamedUniTuple, self).__init__(name)
+        super().__init__(name)
 
     @property
     def iterator_type(self):
@@ -410,7 +410,7 @@ def __init__(self, types, cls):
         self.fields = tuple(cls._fields)
         self.instance_class = cls
         name = "%s(%s)" % (cls.__name__, ", ".join(str(i) for i in self.types))
-        super(NamedTuple, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
@@ -428,7 +428,7 @@ def __init__(self, dtype, reflected=False, initial_value=None):
         self.reflected = reflected
         cls_name = "reflected list" if reflected else "list"
         name = "%s(%s)<iv=%s>" % (cls_name, self.dtype, initial_value)
-        super(List, self).__init__(name=name)
+        super().__init__(name=name)
         InitialValue.__init__(self, initial_value)
 
     def copy(self, dtype=None, reflected=None):
@@ -558,7 +558,7 @@ def __init__(self, dtype, reflected=False):
         self.reflected = reflected
         cls_name = "reflected set" if reflected else "set"
         name = "%s(%s)" % (cls_name, self.dtype)
-        super(Set, self).__init__(name=name)
+        super().__init__(name=name)
 
     @property
     def key(self):
@@ -613,7 +613,7 @@ class SetEntry(Type):
     def __init__(self, set_type):
         self.set_type = set_type
         name = "entry(%s)" % set_type
-        super(SetEntry, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
@@ -638,7 +638,7 @@ def __init__(self, itemty):
             self.__class__.__name__,
             itemty,
         )
-        super(ListType, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
@@ -680,7 +680,7 @@ def __init__(self, parent):
         self.yield_type = self.parent.item_type
         name = "list[{}]".format(self.parent.name)
         iterator_type = ListTypeIteratorType(self)
-        super(ListTypeIterableType, self).__init__(name, iterator_type)
+        super().__init__(name, iterator_type)
 
 
 class ListTypeIteratorType(SimpleIteratorType):
@@ -689,7 +689,7 @@ def __init__(self, iterable):
         self.iterable = iterable
         yield_type = iterable.yield_type
         name = "iter[{}->{}]".format(iterable.parent, yield_type)
-        super(ListTypeIteratorType, self).__init__(name, yield_type)
+        super().__init__(name, yield_type)
 
 
 def _sentry_forbidden_types(key, value):
@@ -721,7 +721,7 @@ def __init__(self, keyty, valty, initial_value=None):
         name = "{}[{},{}]<iv={}>".format(
             self.__class__.__name__, keyty, valty, initial_value
         )
-        super(DictType, self).__init__(name)
+        super().__init__(name)
         InitialValue.__init__(self, initial_value)
 
     def is_precise(self):
@@ -790,7 +790,7 @@ class FakeNamedTuple(pySequence):
         def __init__(self, name, keys):
             self.__name__ = name
             self._fields = tuple(keys)
-            super(LiteralStrKeyDict.FakeNamedTuple, self).__init__()
+            super().__init__()
 
         def __len__(self):
             return len(self._fields)
@@ -855,7 +855,7 @@ def __init__(self, parent):
         name = "items[{}]".format(self.parent.name)
         self.name = name
         iterator_type = DictIteratorType(self)
-        super(DictItemsIterableType, self).__init__(name, iterator_type)
+        super().__init__(name, iterator_type)
 
 
 class DictKeysIterableType(SimpleIterableType):
@@ -868,7 +868,7 @@ def __init__(self, parent):
         name = "keys[{}]".format(self.parent.name)
         self.name = name
         iterator_type = DictIteratorType(self)
-        super(DictKeysIterableType, self).__init__(name, iterator_type)
+        super().__init__(name, iterator_type)
 
 
 class DictValuesIterableType(SimpleIterableType):
@@ -881,7 +881,7 @@ def __init__(self, parent):
         name = "values[{}]".format(self.parent.name)
         self.name = name
         iterator_type = DictIteratorType(self)
-        super(DictValuesIterableType, self).__init__(name, iterator_type)
+        super().__init__(name, iterator_type)
 
 
 class DictIteratorType(SimpleIteratorType):
@@ -892,7 +892,7 @@ def __init__(self, iterable):
         name = "iter[{}->{}],{}".format(
             iterable.parent, yield_type, iterable.name
         )
-        super(DictIteratorType, self).__init__(name, yield_type)
+        super().__init__(name, yield_type)
 
 
 class StructRef(Type):
diff --git a/numba_cuda/numba/cuda/types/cuda_function_type.py b/numba_cuda/numba/cuda/types/cuda_function_type.py
index 93b548545..fb69f5704 100644
--- a/numba_cuda/numba/cuda/types/cuda_function_type.py
+++ b/numba_cuda/numba/cuda/types/cuda_function_type.py
@@ -121,7 +121,7 @@ def __init__(self, nargs, dispatchers):
             types.undefined, (types.undefined,) * nargs, recvr=None
         )
 
-        super(UndefinedFunctionType, self).__init__(signature)
+        super().__init__(signature)
 
         self.dispatchers = dispatchers
 
@@ -159,7 +159,7 @@ def __init__(self, rtype, atypes):
             lst.append(atype.name)
         name = "%s(%s)" % (rtype, ", ".join(lst))
 
-        super(FunctionPrototype, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
diff --git a/numba_cuda/numba/cuda/types/cuda_functions.py b/numba_cuda/numba/cuda/types/cuda_functions.py
index 6d6b21c1d..5605ded81 100644
--- a/numba_cuda/numba/cuda/types/cuda_functions.py
+++ b/numba_cuda/numba/cuda/types/cuda_functions.py
@@ -69,7 +69,7 @@ def argsnkwargs_to_str(args, kwargs):
     return ", ".join(buf)
 
 
-class _ResolutionFailures(object):
+class _ResolutionFailures:
     """Collect and format function resolution failures."""
 
     def __init__(self, context, function_type, args, kwargs, depth=0):
@@ -286,7 +286,7 @@ def __init__(self, template):
         self._impl_keys = {}
         name = "%s(%s)" % (self.__class__.__name__, self.typing_key)
         self._depth = 0
-        super(BaseFunction, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
@@ -390,7 +390,7 @@ def __init__(self, template, this):
             self.typing_key,
             self.this,
         )
-        super(BoundFunction, self).__init__(name)
+        super().__init__(name)
 
     def unify(self, typingctx, other):
         if (
@@ -561,7 +561,7 @@ class Dispatcher(WeakType, Callable, Dummy):
 
     def __init__(self, dispatcher):
         self._store_object(dispatcher)
-        super(Dispatcher, self).__init__("type(%s)" % dispatcher)
+        super().__init__("type(%s)" % dispatcher)
 
     def dump(self, tab=""):
         print(
@@ -673,7 +673,7 @@ def generic(self, args, kws):
             template = GilRequiringDefn
         else:
             template = make_concrete_template("CFuncPtr", sig, [sig])
-        super(ExternalFunctionPointer, self).__init__(template)
+        super().__init__(template)
 
     @property
     def key(self):
@@ -692,7 +692,7 @@ def __init__(self, symbol, sig):
         self.symbol = symbol
         self.sig = sig
         template = typing.make_concrete_template(symbol, symbol, [sig])
-        super(ExternalFunction, self).__init__(template)
+        super().__init__(template)
 
     @property
     def key(self):
@@ -707,7 +707,7 @@ class NamedTupleClass(Callable, Opaque):
     def __init__(self, instance_class):
         self.instance_class = instance_class
         name = "class(%s)" % (instance_class)
-        super(NamedTupleClass, self).__init__(name)
+        super().__init__(name)
 
     def get_call_type(self, context, args, kws):
         # Overridden by the __call__ constructor resolution in
@@ -733,7 +733,7 @@ class NumberClass(Callable, DTypeSpec, Opaque):
     def __init__(self, instance_type):
         self.instance_type = instance_type
         name = "class(%s)" % (instance_type,)
-        super(NumberClass, self).__init__(name)
+        super().__init__(name)
 
     def get_call_type(self, context, args, kws):
         # Overridden by the __call__ constructor resolution in typing.builtins
@@ -768,7 +768,7 @@ def __init__(self, dispatcher_type):
         assert isinstance(dispatcher_type, Dispatcher)
         self.dispatcher_type = dispatcher_type
         name = "recursive(%s)" % (dispatcher_type,)
-        super(RecursiveCall, self).__init__(name)
+        super().__init__(name)
         # Initializing for the first time
         if self._overloads is None:
             self._overloads = {}
diff --git a/numba_cuda/numba/cuda/types/cuda_iterators.py b/numba_cuda/numba/cuda/types/cuda_iterators.py
index ba320bc79..c9bc15e55 100644
--- a/numba_cuda/numba/cuda/types/cuda_iterators.py
+++ b/numba_cuda/numba/cuda/types/cuda_iterators.py
@@ -9,7 +9,7 @@ class RangeType(SimpleIterableType):
     def __init__(self, dtype):
         self.dtype = dtype
         name = "range_state_%s" % (dtype,)
-        super(SimpleIterableType, self).__init__(name)
+        super().__init__(name)
         self._iterator_type = RangeIteratorType(self.dtype)
 
     def unify(self, typingctx, other):
@@ -22,7 +22,7 @@ def unify(self, typingctx, other):
 class RangeIteratorType(SimpleIteratorType):
     def __init__(self, dtype):
         name = "range_iter_%s" % (dtype,)
-        super(SimpleIteratorType, self).__init__(name)
+        super().__init__(name)
         self._yield_type = dtype
 
     def unify(self, typingctx, other):
@@ -50,7 +50,7 @@ def __init__(
             self.arg_types,
             self.has_finalizer,
         )
-        super(Generator, self).__init__(name, yield_type)
+        super().__init__(name, yield_type)
 
     @property
     def key(self):
@@ -75,7 +75,7 @@ def __init__(self, iterable_type):
         self.source_type = iterable_type.iterator_type
         yield_type = Tuple([intp, self.source_type.yield_type])
         name = "enumerate(%s)" % (self.source_type)
-        super(EnumerateType, self).__init__(name, yield_type)
+        super().__init__(name, yield_type)
 
     @property
     def key(self):
@@ -94,7 +94,7 @@ def __init__(self, iterable_types):
         self.source_types = tuple(tp.iterator_type for tp in iterable_types)
         yield_type = Tuple([tp.yield_type for tp in self.source_types])
         name = "zip(%s)" % ", ".join(str(tp) for tp in self.source_types)
-        super(ZipType, self).__init__(name, yield_type)
+        super().__init__(name, yield_type)
 
     @property
     def key(self):
@@ -117,4 +117,4 @@ def __init__(self, array_type):
         else:
             # iteration semantics leads to A order layout
             yield_type = array_type.copy(ndim=array_type.ndim - 1, layout="A")
-        super(ArrayIterator, self).__init__(name, yield_type)
+        super().__init__(name, yield_type)
diff --git a/numba_cuda/numba/cuda/types/cuda_misc.py b/numba_cuda/numba/cuda/types/cuda_misc.py
index 59801da06..a8c466aad 100644
--- a/numba_cuda/numba/cuda/types/cuda_misc.py
+++ b/numba_cuda/numba/cuda/types/cuda_misc.py
@@ -106,7 +106,7 @@ def __init__(self, value):
         # Use helper function to support both hashable and non-hashable
         # values. See discussion in gh #6957.
         self._value_key = get_hashable_key(value)
-        super(Omitted, self).__init__("omitted(default=%r)" % (value,))
+        super().__init__("omitted(default=%r)" % (value,))
 
     @property
     def key(self):
@@ -126,7 +126,7 @@ class VarArg(Type):
 
     def __init__(self, dtype):
         self.dtype = dtype
-        super(VarArg, self).__init__("*%s" % dtype)
+        super().__init__("*%s" % dtype)
 
     @property
     def key(self):
@@ -136,7 +136,7 @@ def key(self):
 class Module(Dummy):
     def __init__(self, pymod):
         self.pymod = pymod
-        super(Module, self).__init__("Module(%s)" % pymod)
+        super().__init__("Module(%s)" % pymod)
 
     @property
     def key(self):
@@ -154,7 +154,7 @@ class MemInfoPointer(Type):
     def __init__(self, dtype):
         self.dtype = dtype
         name = "memory-managed *%s" % dtype
-        super(MemInfoPointer, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
@@ -181,7 +181,7 @@ def __init__(self, dtype, addrspace=None):
             name = "%s_%s*" % (dtype, addrspace)
         else:
             name = "%s*" % dtype
-        super(CPointer, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
@@ -206,7 +206,7 @@ def __init__(self, dtype, count):
         self.dtype = dtype
         self.count = count
         name = "*%s[%d]" % (dtype, count)
-        super(EphemeralArray, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
@@ -220,7 +220,7 @@ class Object(Type):
     def __init__(self, clsobj):
         self.cls = clsobj
         name = "Object(%s)" % clsobj.__name__
-        super(Object, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
@@ -237,7 +237,7 @@ def __init__(self, typ):
         typ = unliteral(typ)
         self.type = typ
         name = "OptionalType(%s)" % self.type
-        super(Optional, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
@@ -303,7 +303,7 @@ def __init__(self, exc_class):
         assert issubclass(exc_class, BaseException)
         name = "%s" % (exc_class.__name__)
         self.exc_class = exc_class
-        super(ExceptionClass, self).__init__(name)
+        super().__init__(name)
 
     def get_call_type(self, context, args, kws):
         return self.get_call_signatures()[0][0]
@@ -332,7 +332,7 @@ def __init__(self, exc_class):
         assert issubclass(exc_class, BaseException)
         name = "%s(...)" % (exc_class.__name__,)
         self.exc_class = exc_class
-        super(ExceptionInstance, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
@@ -344,7 +344,7 @@ def __init__(self, name, members):
         assert members in (2, 3)
         self.members = members
         self.has_step = members >= 3
-        super(SliceType, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
@@ -379,7 +379,7 @@ class ClassInstanceType(Type):
     def __init__(self, class_type):
         self.class_type = class_type
         name = "{0}.{1}".format(self.name_prefix, self.class_type.name)
-        super(ClassInstanceType, self).__init__(name)
+        super().__init__(name)
 
     def get_data_type(self):
         return ClassDataType(self)
@@ -450,7 +450,7 @@ def __init__(
         name = "{0}.{1}#{2:x}<{3}>".format(
             self.name_prefix, self.class_name, id(self), fielddesc
         )
-        super(ClassType, self).__init__(name)
+        super().__init__(name)
 
     def get_call_type(self, context, args, kws):
         return self.ctor_template(context).apply(args, kws)
@@ -491,7 +491,7 @@ class DeferredType(Type):
     def __init__(self):
         self._define = None
         name = "{0}#{1}".format(type(self).__name__, id(self))
-        super(DeferredType, self).__init__(name)
+        super().__init__(name)
 
     def get(self):
         if self._define is None:
@@ -520,7 +520,7 @@ class ClassDataType(Type):
     def __init__(self, classtyp):
         self.class_type = classtyp
         name = "data.{0}".format(self.class_type.name)
-        super(ClassDataType, self).__init__(name)
+        super().__init__(name)
 
 
 class ContextManager(Callable, Phantom):
@@ -530,7 +530,7 @@ class ContextManager(Callable, Phantom):
 
     def __init__(self, cm):
         self.cm = cm
-        super(ContextManager, self).__init__("ContextManager({})".format(cm))
+        super().__init__("ContextManager({})".format(cm))
 
     def get_call_signatures(self):
         if not self.cm.is_callable:
@@ -555,7 +555,7 @@ def get_impl_key(self, sig):
 
 class UnicodeType(IterableType, Hashable):
     def __init__(self, name):
-        super(UnicodeType, self).__init__(name)
+        super().__init__(name)
 
     @property
     def iterator_type(self):
@@ -566,4 +566,4 @@ class UnicodeIteratorType(SimpleIteratorType):
     def __init__(self, dtype):
         name = "iter_unicode"
         self.data = dtype
-        super(UnicodeIteratorType, self).__init__(name, dtype)
+        super().__init__(name, dtype)
diff --git a/numba_cuda/numba/cuda/types/cuda_npytypes.py b/numba_cuda/numba/cuda/types/cuda_npytypes.py
index 71e8e6b66..74b799aa1 100644
--- a/numba_cuda/numba/cuda/types/cuda_npytypes.py
+++ b/numba_cuda/numba/cuda/types/cuda_npytypes.py
@@ -25,7 +25,7 @@ class CharSeq(Type):
     def __init__(self, count):
         self.count = count
         name = "[char x %d]" % count
-        super(CharSeq, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
@@ -46,7 +46,7 @@ class UnicodeCharSeq(Type):
     def __init__(self, count):
         self.count = count
         name = "[unichr x %d]" % count
-        super(UnicodeCharSeq, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
@@ -149,7 +149,7 @@ def __init__(self, fields, size, aligned):
 
         desc = ",".join(descbuf)
         name = "Record({};{};{})".format(desc, self.size, self.aligned)
-        super(Record, self).__init__(name)
+        super().__init__(name)
 
         self.bitwidth = self.dtype.itemsize * 8
 
@@ -276,7 +276,7 @@ def __init__(self, dtype):
         assert isinstance(dtype, Type)
         self._dtype = dtype
         name = "dtype(%s)" % (dtype,)
-        super(DTypeSpec, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
@@ -287,7 +287,7 @@ def dtype(self):
         return self._dtype
 
     def __getitem__(self, arg):
-        res = super(DType, self).__getitem__(arg)
+        res = super().__getitem__(arg)
         return res.copy(dtype=self.dtype)
 
 
@@ -301,7 +301,7 @@ def __init__(self, arrty):
         yield_type = arrty.dtype
         self.dtype = yield_type
         name = "array.flat({arrayty})".format(arrayty=arrty)
-        super(NumpyFlatType, self).__init__(name, yield_type)
+        super().__init__(name, yield_type)
 
     @property
     def key(self):
@@ -319,7 +319,7 @@ def __init__(self, arrty):
         self.array_type = arrty
         yield_type = Tuple((UniTuple(intp, arrty.ndim), arrty.dtype))
         name = "ndenumerate({arrayty})".format(arrayty=arrty)
-        super(NumpyNdEnumerateType, self).__init__(name, yield_type)
+        super().__init__(name, yield_type)
 
     @property
     def key(self):
@@ -346,7 +346,7 @@ def __init__(self, arrays):
         name = "nditer(ndim={ndim}, layout={layout}, inputs={arrays})".format(
             ndim=self.ndim, layout=self.layout, arrays=self.arrays
         )
-        super(NumpyNdIterType, self).__init__(name)
+        super().__init__(name)
 
     @classmethod
     def _compute_layout(cls, arrays):
@@ -445,7 +445,7 @@ def __init__(self, ndim):
         self.ndim = ndim
         yield_type = UniTuple(intp, self.ndim)
         name = "ndindex(ndim={ndim})".format(ndim=ndim)
-        super(NumpyNdIndexType, self).__init__(name, yield_type)
+        super().__init__(name, yield_type)
 
     @property
     def key(self):
@@ -474,7 +474,7 @@ def __init__(
             if not self.aligned:
                 type_name = "unaligned " + type_name
             name = "%s(%s, %sd, %s)" % (type_name, dtype, ndim, layout)
-        super(Array, self).__init__(dtype, ndim, layout, name=name)
+        super().__init__(dtype, ndim, layout, name=name)
 
     @property
     def mangling_args(self):
@@ -572,7 +572,7 @@ def __init__(self, arytype):
         self.dtype = arytype.dtype
         self.ndim = arytype.ndim
         name = "ArrayCTypes(dtype={0}, ndim={1})".format(self.dtype, self.ndim)
-        super(ArrayCTypes, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
@@ -605,7 +605,7 @@ class ArrayFlags(Type):
     def __init__(self, arytype):
         self.array_type = arytype
         name = "ArrayFlags({0})".format(self.array_type)
-        super(ArrayFlags, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
@@ -630,7 +630,7 @@ def __init__(self, dtype, shape):
         self._shape = shape
         name = "nestedarray(%s, %s)" % (dtype, shape)
         ndim = len(shape)
-        super(NestedArray, self).__init__(dtype, ndim, "C", name=name)
+        super().__init__(dtype, ndim, "C", name=name)
 
     @property
     def shape(self):
@@ -666,19 +666,19 @@ def __repr__(self):
 
 class NumPyRandomBitGeneratorType(Type):
     def __init__(self, *args, **kwargs):
-        super(NumPyRandomBitGeneratorType, self).__init__(*args, **kwargs)
+        super().__init__(*args, **kwargs)
         self.name = "NumPyRandomBitGeneratorType"
 
 
 class NumPyRandomGeneratorType(Type):
     def __init__(self, *args, **kwargs):
-        super(NumPyRandomGeneratorType, self).__init__(*args, **kwargs)
+        super().__init__(*args, **kwargs)
         self.name = "NumPyRandomGeneratorType"
 
 
 class PolynomialType(Type):
     def __init__(self, coef, domain=None, window=None, n_args=1):
-        super(PolynomialType, self).__init__(
+        super().__init__(
             name=f"PolynomialType({coef}, {domain}, {domain}, {n_args})"
         )
         self.coef = coef
diff --git a/numba_cuda/numba/cuda/types/cuda_scalars.py b/numba_cuda/numba/cuda/types/cuda_scalars.py
index 77bba9908..328ae5712 100644
--- a/numba_cuda/numba/cuda/types/cuda_scalars.py
+++ b/numba_cuda/numba/cuda/types/cuda_scalars.py
@@ -31,7 +31,7 @@ def parse_integer_signed(name):
 @total_ordering
 class Integer(Number):
     def __init__(self, name, bitwidth=None, signed=None):
-        super(Integer, self).__init__(name)
+        super().__init__(name)
         if bitwidth is None:
             bitwidth = parse_integer_bitwidth(name)
         if signed is None:
@@ -114,7 +114,7 @@ def can_convert_to(self, typingctx, other):
 @total_ordering
 class Float(Number):
     def __init__(self, *args, **kws):
-        super(Float, self).__init__(*args, **kws)
+        super().__init__(*args, **kws)
         # Determine bitwidth
         assert self.name.startswith("float")
         bitwidth = int(self.name[5:])
@@ -132,7 +132,7 @@ def __lt__(self, other):
 @total_ordering
 class Complex(Number):
     def __init__(self, name, underlying_float, **kwargs):
-        super(Complex, self).__init__(name, **kwargs)
+        super().__init__(name, **kwargs)
         self.underlying_float = underlying_float
         # Determine bitwidth
         assert self.name.startswith("complex")
@@ -157,7 +157,7 @@ def __init__(self, unit, *args, **kws):
         name = "%s[%s]" % (self.type_name, unit)
         self.unit = unit
         self.unit_code = npdatetime_helpers.DATETIME_UNITS[self.unit]
-        super(_NPDatetimeBase, self).__init__(name, *args, **kws)
+        super().__init__(name, *args, **kws)
 
     def __lt__(self, other):
         if self.__class__ is not other.__class__:
@@ -202,7 +202,7 @@ def __init__(self, cls, dtype):
             self.dtype,
             self.instance_class.__name__,
         )
-        super(EnumClass, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
@@ -249,7 +249,7 @@ def __init__(self, cls, dtype):
             self.dtype,
             self.instance_class.__name__,
         )
-        super(EnumMember, self).__init__(name)
+        super().__init__(name)
 
     @property
     def key(self):
diff --git a/numba_cuda/numba/cuda/typing/builtins.py b/numba_cuda/numba/cuda/typing/builtins.py
index c380cbf43..7ffb34c99 100644
--- a/numba_cuda/numba/cuda/typing/builtins.py
+++ b/numba_cuda/numba/cuda/typing/builtins.py
@@ -896,7 +896,7 @@ def resolve___call__(self, classty):
             # For example, see numba/typed/typeddict.py
             #   @type_callable(DictType)
             #   def typeddict_call(context):
-            class Redirect(object):
+            class Redirect:
                 def __init__(self, context):
                     self.context = context
 
@@ -1159,7 +1159,7 @@ def generic_resolve(self, deferred, attr):
 # ------------------------------------------------------------------------------
 
 
-class IndexValue(object):
+class IndexValue:
     """
     Index and value
     """
@@ -1175,9 +1175,7 @@ def __repr__(self):
 class IndexValueType(types.Type):
     def __init__(self, val_typ):
         self.val_typ = val_typ
-        super(IndexValueType, self).__init__(
-            name="IndexValueType({})".format(val_typ)
-        )
+        super().__init__(name="IndexValueType({})".format(val_typ))
 
 
 @typeof_impl.register(IndexValue)
diff --git a/numba_cuda/numba/cuda/typing/context.py b/numba_cuda/numba/cuda/typing/context.py
index 9514d1e8b..c9d4c6832 100644
--- a/numba_cuda/numba/cuda/typing/context.py
+++ b/numba_cuda/numba/cuda/typing/context.py
@@ -18,7 +18,7 @@
 from numba.cuda.utils import order_by_target_specificity
 
 
-class Rating(object):
+class Rating:
     __slots__ = "promote", "safe_convert", "unsafe_convert"
 
     def __init__(self):
@@ -103,7 +103,7 @@ def match(self, py_func, args):
                 return frame
 
 
-class CallFrame(object):
+class CallFrame:
     """
     A compile-time call frame
     """
@@ -131,7 +131,7 @@ def add_return_type(self, return_type):
             raise errors.TypingError(m)
 
 
-class BaseContext(object):
+class BaseContext:
     """A typing context for storing function typing constrain template."""
 
     def __init__(self):
diff --git a/numba_cuda/numba/cuda/typing/npydecl.py b/numba_cuda/numba/cuda/typing/npydecl.py
index 93eab7a44..029383286 100644
--- a/numba_cuda/numba/cuda/typing/npydecl.py
+++ b/numba_cuda/numba/cuda/typing/npydecl.py
@@ -238,7 +238,7 @@ def generic(self, args, kws):
         (particularly user-defined operators).
         """
         try:
-            sig = super(NumpyRulesArrayOperator, self).generic(args, kws)
+            sig = super().generic(args, kws)
         except TypingError:
             return None
         if sig is None:
@@ -279,7 +279,7 @@ def generic(self, args, kws):
         if not isinstance(lhs, types.ArrayCompatible):
             return
         args = args + (lhs,)
-        sig = super(NumpyRulesInplaceArrayOperator, self).generic(args, kws)
+        sig = super().generic(args, kws)
         # Strip off the fake explicit output
         assert len(sig.args) == 3
         real_sig = signature(sig.return_type, *sig.args[:2])
@@ -296,7 +296,7 @@ class NumpyRulesUnaryArrayOperator(NumpyRulesArrayOperator):
     def generic(self, args, kws):
         assert not kws
         if len(args) == 1 and isinstance(args[0], types.ArrayCompatible):
-            return super(NumpyRulesUnaryArrayOperator, self).generic(args, kws)
+            return super().generic(args, kws)
 
 
 # list of unary ufuncs to register
diff --git a/numba_cuda/numba/cuda/typing/templates.py b/numba_cuda/numba/cuda/typing/templates.py
index f9d83f95a..42dda278c 100644
--- a/numba_cuda/numba/cuda/typing/templates.py
+++ b/numba_cuda/numba/cuda/typing/templates.py
@@ -34,7 +34,7 @@
 _inline_info = namedtuple("inline_info", "func_ir typemap calltypes signature")
 
 
-class Signature(object):
+class Signature:
     """
     The signature of a function call or operation, i.e. its argument types
     and return type.
@@ -553,7 +553,7 @@ def get_template_info(self):
 
 class _EmptyImplementationEntry(InternalError):
     def __init__(self, reason):
-        super(_EmptyImplementationEntry, self).__init__(
+        super().__init__(
             "_EmptyImplementationEntry({!r})".format(reason),
         )
 
@@ -956,7 +956,7 @@ def make_overload_template(
     return type(base)(name, (base,), dct)
 
 
-class _TemplateTargetHelperMixin(object):
+class _TemplateTargetHelperMixin:
     """Mixin for helper methods that assist with target/registry resolution"""
 
     def _get_target_registry(self, reason):
@@ -1105,7 +1105,7 @@ def make_intrinsic_template(
     return type(base)(name, (base,), dct)
 
 
-class AttributeTemplate(object):
+class AttributeTemplate:
     def __init__(self, context):
         self.context = context
 
@@ -1137,7 +1137,7 @@ class _OverloadAttributeTemplate(_TemplateTargetHelperMixin, AttributeTemplate):
     is_method = False
 
     def __init__(self, context):
-        super(_OverloadAttributeTemplate, self).__init__(context)
+        super().__init__(context)
         self.context = context
         self._init_once()
 
@@ -1339,7 +1339,7 @@ def generic(_, args, kws):
 # -----------------------------
 
 
-class Registry(object):
+class Registry:
     """
     A registry of typing declarations.  The registry stores such declarations
     for functions, attributes and globals.
@@ -1406,7 +1406,7 @@ def decorator(cls):
             return decorator
 
 
-class BaseRegistryLoader(object):
+class BaseRegistryLoader:
     """
     An incremental loader for a registry.  Each new call to
     new_registrations() will iterate over the not yet seen registrations.
diff --git a/numba_cuda/numba/cuda/utils.py b/numba_cuda/numba/cuda/utils.py
index da20c661a..23f5956d6 100644
--- a/numba_cuda/numba/cuda/utils.py
+++ b/numba_cuda/numba/cuda/utils.py
@@ -261,7 +261,7 @@ def enter(self, state):
             self.pop()
 
 
-class ConfigOptions(object):
+class ConfigOptions:
     OPTIONS = {}
 
     def __init__(self):
@@ -285,7 +285,7 @@ def __getattr__(self, name):
 
     def __setattr__(self, name, value):
         if name.startswith("_"):
-            super(ConfigOptions, self).__setattr__(name, value)
+            super().__setattr__(name, value)
         else:
             self._check_attr(name)
             self._values[name] = value
@@ -359,7 +359,7 @@ class UniqueDict(dict):
     def __setitem__(self, key, value):
         if key in self:
             raise AssertionError("key already in dictionary: %r" % (key,))
-        super(UniqueDict, self).__setitem__(key, value)
+        super().__setitem__(key, value)
 
 
 def runonce(fn):
@@ -406,7 +406,7 @@ def sublist_iterator(start, stop):
         start = stop
 
 
-class BenchmarkResult(object):
+class BenchmarkResult:
     def __init__(self, func, records, loop):
         self.func = func
         self.loop = loop
@@ -672,7 +672,7 @@ def dump_llvm(fndesc, module):
     print("=" * 80)
 
 
-class _lazy_pformat(object):
+class _lazy_pformat:
     def __init__(self, *args, **kwargs):
         self.args = args
         self.kwargs = kwargs
diff --git a/numba_cuda/numba/cuda/vectorizers.py b/numba_cuda/numba/cuda/vectorizers.py
index 2607ea043..5bc53c335 100644
--- a/numba_cuda/numba/cuda/vectorizers.py
+++ b/numba_cuda/numba/cuda/vectorizers.py
@@ -11,7 +11,7 @@
 )
 
 
-class CUDAUFuncDispatcher(object):
+class CUDAUFuncDispatcher:
     """
     Invoke the CUDA ufunc specialization for the given inputs.
     """