Skip to content

Commit a7f5393

Browse files
jiel-nvjiel-nv
committed
another fix for ruff format
1 parent 98e1283 commit a7f5393

File tree

2 files changed

+31
-33
lines changed

2 files changed

+31
-33
lines changed

‎numba_cuda/numba/cuda/debuginfo.py‎

Lines changed: 23 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -18,15 +18,17 @@ def _get_llvmlite_version():
1818
"""Get llvmlite version as tuple (major, minor, patch)."""
1919
try:
2020
import llvmlite
21+
2122
version_str = llvmlite.__version__
2223
# Parse version string like "0.46.0" or "0.46.0dev"
23-
parts = version_str.split('.')
24+
parts = version_str.split(".")
2425
major = int(parts[0])
2526
minor = int(parts[1])
2627
return (major, minor)
2728
except Exception:
2829
return (0, 0)
2930

31+
3032
def _check_polymorphic_debug_info_support():
3133
"""Check if CTK and llvmlite support polymorphic debug info.
3234
@@ -38,6 +40,7 @@ def _check_polymorphic_debug_info_support():
3840
"""
3941
try:
4042
from numba.cuda.cudadrv import runtime
43+
4144
ctk_version = runtime.get_version()
4245
llvmlite_version = _get_llvmlite_version()
4346

@@ -53,17 +56,18 @@ def _check_polymorphic_debug_info_support():
5356
except Exception:
5457
return (False, False)
5558

59+
5660
# Check support and determine mode
57-
(DEBUG_POLY_SUPPORTED,
58-
DEBUG_POLY_USE_TYPED_CONST) = _check_polymorphic_debug_info_support()
61+
(DEBUG_POLY_SUPPORTED, DEBUG_POLY_USE_TYPED_CONST) = (
62+
_check_polymorphic_debug_info_support()
63+
)
5964

6065
# Set config based on polymorphic debug info support
6166
if not hasattr(config, "CUDA_DEBUG_POLY"):
6267
config.CUDA_DEBUG_POLY = DEBUG_POLY_SUPPORTED
6368
if not hasattr(config, "CUDA_DEBUG_POLY_USE_TYPED_CONST"):
64-
config.CUDA_DEBUG_POLY_USE_TYPED_CONST = (
65-
DEBUG_POLY_USE_TYPED_CONST
66-
)
69+
config.CUDA_DEBUG_POLY_USE_TYPED_CONST = DEBUG_POLY_USE_TYPED_CONST
70+
6771

6872
@contextmanager
6973
def suspend_emission(builder):
@@ -699,13 +703,13 @@ def _var_type(self, lltype, size, datamodel=None):
699703
# Polymorphic debug info with DW_TAG_variant
700704
# extraData depends on llvmlite version
701705
if config.CUDA_DEBUG_POLY_USE_TYPED_CONST:
702-
metadata_dict["extraData"] = (
703-
ir.IntType(8)(index)
706+
metadata_dict["extraData"] = ir.IntType(8)(
707+
index
704708
)
705709
else:
706710
# Use metadata node reference
707-
metadata_dict["extraData"] = (
708-
m.add_metadata([ir.IntType(8)(index)])
711+
metadata_dict["extraData"] = m.add_metadata(
712+
[ir.IntType(8)(index)]
709713
)
710714
# Add offset to each variant member
711715
# Offset equals the element's own width
@@ -729,11 +733,14 @@ def _var_type(self, lltype, size, datamodel=None):
729733
{
730734
"tag": ir.DIToken("DW_TAG_member"),
731735
"name": "discriminator",
732-
"baseType": m.add_debug_info("DIBasicType", {
733-
"name": "int",
734-
"size": _BYTE_SIZE,
735-
"encoding": ir.DIToken("DW_ATE_unsigned")
736-
}),
736+
"baseType": m.add_debug_info(
737+
"DIBasicType",
738+
{
739+
"name": "int",
740+
"size": _BYTE_SIZE,
741+
"encoding": ir.DIToken("DW_ATE_unsigned"),
742+
},
743+
),
737744
"size": _BYTE_SIZE,
738745
"flags": ir.DIToken("DIFlagArtificial"),
739746
},
@@ -782,7 +789,7 @@ def _var_type(self, lltype, size, datamodel=None):
782789
"elements": m.add_metadata(meta),
783790
"size": maxwidth,
784791
},
785-
is_distinct = True,
792+
is_distinct=True,
786793
)
787794
# For other cases, use upstream Numba implementation
788795
return super()._var_type(lltype, size, datamodel=datamodel)

‎numba_cuda/numba/cuda/lowering.py‎

Lines changed: 8 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -1709,31 +1709,26 @@ def storevar(self, value, name, argidx=None):
17091709
# Bitcast union pointer directly to i8* and write
17101710
# discriminant at offset 0
17111711
discriminant_ptr = self.builder.bitcast(
1712-
ptr,
1713-
llvm_ir.PointerType(llvm_ir.IntType(8))
1712+
ptr, llvm_ir.PointerType(llvm_ir.IntType(8))
17141713
)
17151714
discriminant_i8 = llvm_ir.Constant(
1716-
llvm_ir.IntType(8),
1717-
discriminant_val
1715+
llvm_ir.IntType(8), discriminant_val
17181716
)
17191717
self.builder.store(discriminant_i8, discriminant_ptr)
17201718
# Secondly write data at offset = sizeof(fetype) in bytes
17211719
lltype = self.context.get_value_type(fetype)
17221720
sizeof_bytes = self.context.get_abi_sizeof(lltype)
17231721
# Bitcast to i8* and use byte-level GEP
17241722
byte_ptr = self.builder.bitcast(
1725-
ptr,
1726-
llvm_ir.PointerType(llvm_ir.IntType(8))
1723+
ptr, llvm_ir.PointerType(llvm_ir.IntType(8))
17271724
)
17281725
data_byte_ptr = self.builder.gep(
17291726
byte_ptr,
1730-
[llvm_ir.Constant(llvm_ir.IntType(32),
1731-
sizeof_bytes)],
1727+
[llvm_ir.Constant(llvm_ir.IntType(32), sizeof_bytes)],
17321728
)
17331729
# Cast to the correct type pointer
17341730
castptr = self.builder.bitcast(
1735-
data_byte_ptr,
1736-
llvm_ir.PointerType(lltype)
1731+
data_byte_ptr, llvm_ir.PointerType(lltype)
17371732
)
17381733
self.builder.store(value, castptr)
17391734
return
@@ -1928,23 +1923,19 @@ def getvar(self, name):
19281923
sizeof_bytes = self.context.get_abi_sizeof(lltype)
19291924
# Bitcast to i8* and use byte-level GEP
19301925
byte_ptr = self.builder.bitcast(
1931-
ptr,
1932-
llvm_ir.PointerType(llvm_ir.IntType(8))
1926+
ptr, llvm_ir.PointerType(llvm_ir.IntType(8))
19331927
)
19341928
value_byte_ptr = self.builder.gep(
19351929
byte_ptr,
1936-
[llvm_ir.Constant(llvm_ir.IntType(32),
1937-
sizeof_bytes)],
1930+
[llvm_ir.Constant(llvm_ir.IntType(32), sizeof_bytes)],
19381931
)
19391932
# Cast to the correct type pointer
19401933
castptr = self.builder.bitcast(
19411934
value_byte_ptr, llvm_ir.PointerType(lltype)
19421935
)
19431936
else:
19441937
# Otherwise, just bitcast to the correct type
1945-
castptr = self.builder.bitcast(
1946-
ptr, llvm_ir.PointerType(lltype)
1947-
)
1938+
castptr = self.builder.bitcast(ptr, llvm_ir.PointerType(lltype))
19481939
return castptr
19491940
else:
19501941
return super().getvar(name)

0 commit comments

Comments
 (0)