diff --git a/transformer_engine/__init__.py b/transformer_engine/__init__.py index e51f03e3d8..c3fb004659 100644 --- a/transformer_engine/__init__.py +++ b/transformer_engine/__init__.py @@ -10,6 +10,37 @@ from importlib import metadata import transformer_engine.common +import torch + +# Public, simple global (kept for backward compatibility). +TE_DEVICE_TYPE = "cuda" +TE_PLATFORM = torch.cuda + +# Apply MUSA (VENDOR) Patches, such as torch.cuda.device -> torch.musa.device +try: + from .plugin.core.backends.vendor.musa.patches import apply_patch as _musa_apply_patch + + _musa_apply_patch() + print("[TE-FL] MUSA patches applied") +except Exception as e: + print(f"[TE-FL] MUSA patches not applied: {e}") + pass + + +def te_device_type(default: str = "cuda") -> str: + try: + return TE_DEVICE_TYPE + except Exception: + return default + + +def te_platform(default=torch.cuda): + try: + return TE_PLATFORM + except Exception: + return default + + try: from . import pytorch except ImportError: diff --git a/transformer_engine/debug/features/fake_quant.py b/transformer_engine/debug/features/fake_quant.py index 58c7379b5b..00c1096351 100644 --- a/transformer_engine/debug/features/fake_quant.py +++ b/transformer_engine/debug/features/fake_quant.py @@ -14,6 +14,7 @@ import transformer_engine_torch as tex +from transformer_engine import te_device_type from transformer_engine.debug.features.api import TEConfigAPIMapper from transformer_engine.common.recipe import Format from transformer_engine.pytorch.tensor import Quantizer @@ -30,7 +31,9 @@ def fake_quantize(tensor: torch.Tensor, fp8_format: tex.DType, out=None): torch.float16, torch.bfloat16, ), "[NVTORCH INSPECT ERROR] Unsupported tensor type." - assert tensor.is_cuda, "[NVTORCH INSPECT ERROR] Must be a GPU tensor." + assert ( + tensor.device.type == te_device_type() + ), f"[NVTORCH INSPECT ERROR] Must be a {te_device_type()} tensor." assert fp8_format in { "FP8E4M3", "FP8E5M2", diff --git a/transformer_engine/debug/features/log_fp8_tensor_stats.py b/transformer_engine/debug/features/log_fp8_tensor_stats.py index d09fb10579..290eb8c35d 100644 --- a/transformer_engine/debug/features/log_fp8_tensor_stats.py +++ b/transformer_engine/debug/features/log_fp8_tensor_stats.py @@ -14,6 +14,7 @@ from nvdlfw_inspect.debug_features.log_tensor_stats import LogTensorStats as BaseLogTensorStats from nvdlfw_inspect.registry import Registry, api_method +from transformer_engine import te_device_type from transformer_engine.debug.features.utils.stats_buffer import STATS_BUFFERS from transformer_engine.pytorch.tensor import Quantizer, QuantizedTensor from transformer_engine.pytorch.tensor.float8_tensor import ( @@ -47,7 +48,10 @@ def _get_new_quantizer(recipe_name, fp8_dtype): return Float8BlockQuantizer(fp8_dtype=fp8_dtype, rowwise=True, columnwise=True) if recipe_name == "fp8_current_scaling": return Float8CurrentScalingQuantizer( - fp8_dtype=fp8_dtype, device=torch.device("cuda"), rowwise=True, columnwise=True + fp8_dtype=fp8_dtype, + device=torch.device(te_device_type()), + rowwise=True, + columnwise=True, ) if recipe_name == "mxfp8": return MXFP8Quantizer(fp8_dtype=fp8_dtype, rowwise=True, columnwise=True) diff --git a/transformer_engine/debug/features/per_tensor_scaling.py b/transformer_engine/debug/features/per_tensor_scaling.py index dd1f42cf06..10ee77a474 100644 --- a/transformer_engine/debug/features/per_tensor_scaling.py +++ b/transformer_engine/debug/features/per_tensor_scaling.py @@ -11,7 +11,9 @@ import nvdlfw_inspect.api as debug_api from nvdlfw_inspect.registry import Registry, api_method + import transformer_engine_torch as tex +from transformer_engine import te_device_type from transformer_engine.pytorch.tensor import Quantizer from transformer_engine.pytorch.tensor.float8_tensor import ( Float8Tensor, @@ -33,7 +35,9 @@ def per_tensor_cast( torch.float16, torch.bfloat16, ), "[NVTORCH INSPECT ERROR] Unsupported tensor type for per tensor current scaling" - assert tensor.is_cuda, "[NVTORCH INSPECT ERROR] Must be a GPU tensor." + assert ( + tensor.device.type == te_device_type() + ), f"[NVTORCH INSPECT ERROR] Must be a {te_device_type()} tensor." assert fp8_dtype in { tex.DType.kFloat8E4M3, tex.DType.kFloat8E5M2, diff --git a/transformer_engine/debug/features/utils/stats_buffer.py b/transformer_engine/debug/features/utils/stats_buffer.py index 20236fb950..e570443d5b 100644 --- a/transformer_engine/debug/features/utils/stats_buffer.py +++ b/transformer_engine/debug/features/utils/stats_buffer.py @@ -16,6 +16,7 @@ from nvdlfw_inspect.utils import gather_along_first_dim from nvdlfw_inspect.logging import MetricLogger +from transformer_engine import te_device_type from transformer_engine.debug.features.utils.stats_computation import ( STATS, DEPENDENCIES, @@ -41,14 +42,14 @@ def __init__(self, layer_name, tensor_name, stats, reduction_group, reduce_withi for stat in stats: self.stats_to_compute = self.stats_to_compute | DEPENDENCIES[stat] - self._buffer = torch.zeros(len(STATS), dtype=torch.float32).cuda() + self._buffer = torch.zeros(len(STATS), dtype=torch.float32).to(te_device_type()) self._new_buffer = self._buffer.clone() self._tmp_buffer = self._buffer.clone() # in case of data parallelism it is possible that layer will not be run on one node # modified is set to True if node is run # we do not take not run nodes into account - self.modified = torch.tensor([False], dtype=torch.bool).cuda() + self.modified = torch.tensor([False], dtype=torch.bool).to(te_device_type()) self.iteration = None self.skip_reduction = False diff --git a/transformer_engine/plugin/core/backends/flagos/attention/dot_product_attention/backends.py b/transformer_engine/plugin/core/backends/flagos/attention/dot_product_attention/backends.py index 8f2e9aeb41..f967dc54d8 100644 --- a/transformer_engine/plugin/core/backends/flagos/attention/dot_product_attention/backends.py +++ b/transformer_engine/plugin/core/backends/flagos/attention/dot_product_attention/backends.py @@ -9,6 +9,7 @@ from packaging.version import Version as PkgVersion import torch +from transformer_engine import te_device_type from transformer_engine.pytorch.utils import ( get_device_compute_capability, ) @@ -283,8 +284,10 @@ def _forward_impl( for x in [query_layer, key_layer, value_layer] ), "FLAttention only supports FP16 and BF16 data types, or Float8Tensors." assert ( - query_layer.is_cuda and key_layer.is_cuda and value_layer.is_cuda - ), "FLAttention only supports CUDA tensors." + query_layer.device.type == te_device_type() + and key_layer.device.type == te_device_type() + and value_layer.device.type == te_device_type() + ), f"FLAttention only supports {te_device_type()} tensors." assert qkv_layout in QKVLayouts, f"FLAttention does not support qkv_layout = {qkv_layout}!" cp_size = 1 diff --git a/transformer_engine/plugin/core/backends/flagos/flagos.py b/transformer_engine/plugin/core/backends/flagos/flagos.py index fd8a61f492..d33bcf1411 100644 --- a/transformer_engine/plugin/core/backends/flagos/flagos.py +++ b/transformer_engine/plugin/core/backends/flagos/flagos.py @@ -243,7 +243,7 @@ def get_cudnn_version(self) -> int: return 90000 def get_num_cublas_streams(self) -> int: - return 0 + return 4 # keep consistent with transformer_engine/common/util/multi_stream.cpp, get_num_compute_streams() ############## class func ################################# def get_flash_attention_class(self): diff --git a/transformer_engine/plugin/core/backends/flagos/register_ops.py b/transformer_engine/plugin/core/backends/flagos/register_ops.py index 0136b6a983..d744cdda41 100644 --- a/transformer_engine/plugin/core/backends/flagos/register_ops.py +++ b/transformer_engine/plugin/core/backends/flagos/register_ops.py @@ -124,6 +124,22 @@ def register_builtins(registry) -> None: vendor=None, priority=150, ), + OpImpl( + op_name="get_num_cublas_streams", + impl_id="default.flagos", + kind=BackendImplKind.DEFAULT, + fn=_bind_is_available(backend.get_num_cublas_streams, is_avail), + vendor=None, + priority=150, + ), + OpImpl( + op_name="get_cudnn_version", + impl_id="default.flagos", + kind=BackendImplKind.DEFAULT, + fn=_bind_is_available(backend.get_cudnn_version, is_avail), + vendor=None, + priority=150, + ), ] registry.register_many(impls) diff --git a/transformer_engine/plugin/core/backends/reference/impl/softmax.py b/transformer_engine/plugin/core/backends/reference/impl/softmax.py index 1783ada92b..2689ab938a 100644 --- a/transformer_engine/plugin/core/backends/reference/impl/softmax.py +++ b/transformer_engine/plugin/core/backends/reference/impl/softmax.py @@ -44,20 +44,35 @@ def scaled_masked_softmax_forward_torch( mask: torch.Tensor, scale: float, ) -> torch.Tensor: - # Handle uint8 mask (CUDA format: 1=masked, 0=unmasked) - # Convert to additive mask (-10000 for masked positions, 0 for unmasked) - if mask.dtype == torch.uint8: - additive_mask = torch.zeros_like(input, dtype=input.dtype) - # Expand mask if needed (mask shape: batch, 1, seq_q, seq_k) - if mask.dim() == 4 and mask.size(1) == 1 and input.dim() == 4: - mask = mask.expand_as(input) - additive_mask = additive_mask.masked_fill(mask.bool(), -10000.0) - else: - additive_mask = mask - - scaled_input = input * scale + additive_mask - - return F.softmax(scaled_input, dim=-1) + """Reference forward matching TE CUDA `scaled_masked_softmax_warp_forward`. + + Integer/bool mask (same as uint8 kernel contract): + - **Exactly** ``mask == 1`` means **masked** (logit set to ``-10000``, not ``input*scale`` offset). + - Any other value (typically 0) means **unmasked** (logit is ``input * scale``). + + Floating mask: treated as **additive** bias in logit space (already scaled), added after + ``input * scale``. + + Common pitfalls this avoids vs the old implementation: + 1) ``input * scale + (-10000)`` on masked positions ≠ CUDA's plain ``-10000``. + 2) Non-uint8 masks (bool, int) were used as direct addends → wrong (0/1 added to logits). + 3) ``mask.bool()`` masks any nonzero byte; CUDA only masks when ``mask == 1``. + """ + if mask.dim() == 4 and mask.size(1) == 1 and input.dim() == 4: + mask = mask.expand_as(input) + + scaled = input * scale + + if mask.is_floating_point(): + scaled = scaled + mask.to(dtype=scaled.dtype) + return F.softmax(scaled, dim=-1) + + # Integer / bool: align with CUDA (masked iff value == 1) + scaled = scaled.masked_fill(mask == 1, -10000.0) + # CUDA zeros output row when every position in the softmax dim is masked (max == -10000) + all_masked = (mask == 1).all(dim=-1, keepdim=True) + out = F.softmax(scaled, dim=-1) + return out.masked_fill(all_masked, 0.0) def scaled_masked_softmax_backward_torch( diff --git a/transformer_engine/plugin/core/backends/reference/reference.py b/transformer_engine/plugin/core/backends/reference/reference.py index 984d62022f..9755d85373 100644 --- a/transformer_engine/plugin/core/backends/reference/reference.py +++ b/transformer_engine/plugin/core/backends/reference/reference.py @@ -476,7 +476,7 @@ def get_cudnn_version(self) -> int: return 0 def get_num_cublas_streams(self) -> int: - return 0 + return 4 # keep consistent with transformer_engine/common/util/multi_stream.cpp, get_num_compute_streams() # Multi-tensor functions def multi_tensor_scale( diff --git a/transformer_engine/plugin/core/backends/vendor/musa/patches.py b/transformer_engine/plugin/core/backends/vendor/musa/patches.py new file mode 100644 index 0000000000..220c5be03f --- /dev/null +++ b/transformer_engine/plugin/core/backends/vendor/musa/patches.py @@ -0,0 +1,72 @@ +"""Python-side compatibility patches for the MUSA vendor backend.""" + +from __future__ import annotations + +from collections.abc import Callable + +import torch + + +def _noop(*args, **kwargs): + return None + + +# Patches: (parent_object, attribute_name, replacement_callable) +_PATCH_CALLS: list[tuple[object, str, Callable[..., object]]] = [ + # We do not recommend replace is_available, due to its device-related behavior. + # (torch.cuda, "is_available", torch.musa.is_available), + (torch.cuda, "get_device_properties", torch.musa.get_device_properties), + (torch.cuda, "device", torch.musa.device), + (torch.cuda, "current_device", torch.musa.current_device), + (torch.cuda, "synchronize", torch.musa.synchronize), + (torch.cuda, "is_current_stream_capturing", torch.musa.is_current_stream_capturing), + # TODO: Add NVTX patches for MUSA. + # NVTX is CUDA-specific; make it a no-op on MUSA. + (torch.cuda.nvtx, "range_push", _noop), + (torch.cuda.nvtx, "range_pop", _noop), + # TODO: Add other patches for MUSA. +] + + +def apply_patch() -> None: + """Apply MUSA Python-side patches (idempotent, best-effort).""" + try: + from .musa import MUSABackend + + if not MUSABackend().is_available(): + return + except Exception as e: + print(f"[TE-FL] MUSA backend not available: {e}") + # If backend availability can't be determined, don't patch. + return + + # Mark TE global device type for Python-side callers. + # IMPORTANT: do not import `transformer_engine` here, because TE's `__init__.py` + # imports this module to run patches and that would cause a circular import. + try: + import transformer_engine + + transformer_engine.TE_DEVICE_TYPE = "musa" + transformer_engine.TE_PLATFORM = torch.musa + except Exception as e: + print(f"[TE-FL Musa Patches] Error setting TE device type or platform: {e}") + # Best-effort: don't fail patching if we can't set the global. + pass + + # Only patch when torch.musa exists and is usable. + if not hasattr(torch, "musa"): + return + try: + if not torch.musa.is_available(): + return + except Exception: + return + + for parent, attr, replacement in _PATCH_CALLS: + if not hasattr(parent, attr): + continue + try: + setattr(parent, attr, replacement) + except Exception: + # Best-effort: patching should never crash import/initialization. + continue diff --git a/transformer_engine/pytorch/attention/dot_product_attention/backends.py b/transformer_engine/pytorch/attention/dot_product_attention/backends.py index 95558e30da..270e6a2ee8 100644 --- a/transformer_engine/pytorch/attention/dot_product_attention/backends.py +++ b/transformer_engine/pytorch/attention/dot_product_attention/backends.py @@ -15,6 +15,7 @@ import torch import torch.nn.functional as F import transformer_engine_torch as tex +from transformer_engine import te_device_type from transformer_engine.pytorch.utils import ( get_device_compute_capability, split_tensor_along_dim, @@ -387,10 +388,10 @@ def forward( fp8_recipe = fp8_meta["local_recipes"][0] if fp8_recipe.float8_current_scaling(): S_quantizer = Float8CurrentScalingQuantizer( - fp8_dtype=S_quantizer.dtype, device="cuda" + fp8_dtype=S_quantizer.dtype, device=te_device_type() ) dP_quantizer = Float8CurrentScalingQuantizer( - fp8_dtype=dP_quantizer.dtype, device="cuda" + fp8_dtype=dP_quantizer.dtype, device=te_device_type() ) if "2" in qkv_layout or "3" in qkv_layout: @@ -676,8 +677,10 @@ def forward( for x in [query_layer, key_layer, value_layer] ), "FlashAttention only supports FP16 and BF16 data types, or Float8Tensors." assert ( - query_layer.is_cuda and key_layer.is_cuda and value_layer.is_cuda - ), "FlashAttention currently only supports CUDA tensors." + query_layer.device.type == te_device_type() + and key_layer.device.type == te_device_type() + and value_layer.device.type == te_device_type() + ), f"FlashAttention currently only supports {te_device_type()} tensors." assert ( qkv_layout in QKVLayouts ), f"FlashAttention does not support qkv_layout = {qkv_layout}!" @@ -1738,8 +1741,10 @@ def forward( for x in [query_layer, key_layer, value_layer] ), "FusedAttention only supports FP16 and BF16 data types, or Float8Tensors." assert ( - query_layer.is_cuda and key_layer.is_cuda and value_layer.is_cuda - ), "FusedAttention only supports CUDA tensors." + query_layer.device.type == te_device_type() + and key_layer.device.type == te_device_type() + and value_layer.device.type == te_device_type() + ), f"FusedAttention only supports {te_device_type()} tensors." assert ( qkv_layout in QKVLayouts ), f"FusedAttention does not support qkv_layout = {qkv_layout}!" diff --git a/transformer_engine/pytorch/attention/dot_product_attention/dot_product_attention.py b/transformer_engine/pytorch/attention/dot_product_attention/dot_product_attention.py index 4e5a79e668..8c96f66aaa 100644 --- a/transformer_engine/pytorch/attention/dot_product_attention/dot_product_attention.py +++ b/transformer_engine/pytorch/attention/dot_product_attention/dot_product_attention.py @@ -14,6 +14,7 @@ from torch.nn.parameter import Parameter import transformer_engine_torch as tex +from transformer_engine import te_device_type from transformer_engine.common.recipe import ( Format, Recipe, @@ -420,12 +421,14 @@ def __init__( self.softmax_offset = None if self.softmax_type == "off-by-one": self.softmax_offset = torch.zeros( - self.num_attention_heads // self.tp_size, device="cuda" + self.num_attention_heads // self.tp_size, device=te_device_type() ) if self.softmax_type == "learnable": self.register_parameter( "softmax_offset", - Parameter(torch.empty(self.num_attention_heads // self.tp_size, device="cuda")), + Parameter( + torch.empty(self.num_attention_heads // self.tp_size, device=te_device_type()) + ), get_rng_state_tracker=get_rng_state_tracker, ) @@ -1026,8 +1029,10 @@ def forward( # checks for q/k/v shapes assert ( - query_layer.is_cuda and key_layer.is_cuda and value_layer.is_cuda - ), "DotProductAttention only supports CUDA tensors." + query_layer.device.type == te_device_type() + and key_layer.device.type == te_device_type() + and value_layer.device.type == te_device_type() + ), f"DotProductAttention only supports {te_device_type()} tensors." assert ( query_layer.dtype == key_layer.dtype and query_layer.dtype == value_layer.dtype ), "Queries, keys and values must have the same data type!" diff --git a/transformer_engine/pytorch/attention/dot_product_attention/softmax.py b/transformer_engine/pytorch/attention/dot_product_attention/softmax.py index df10fc7905..57e5d4f425 100644 --- a/transformer_engine/pytorch/attention/dot_product_attention/softmax.py +++ b/transformer_engine/pytorch/attention/dot_product_attention/softmax.py @@ -8,6 +8,7 @@ import torch from torch import nn import transformer_engine_torch as tex +from transformer_engine import te_device_type from transformer_engine.pytorch.export import is_in_onnx_export_mode @@ -24,7 +25,7 @@ def _get_default_causal_mask(mask_type: str, sq: int, sk: int) -> torch.Tensor: def _get_mask(): diagonal_offset = sk - sq + 1 if "bottom_right" in mask_type else 1 return torch.triu( - torch.ones(sq, sk, dtype=torch.bool, device="cuda"), diagonal=diagonal_offset + torch.ones(sq, sk, dtype=torch.bool, device=te_device_type()), diagonal=diagonal_offset ) if is_in_onnx_export_mode(): diff --git a/transformer_engine/pytorch/attention/dot_product_attention/utils.py b/transformer_engine/pytorch/attention/dot_product_attention/utils.py index 6bcc9f25da..ae36eb4160 100644 --- a/transformer_engine/pytorch/attention/dot_product_attention/utils.py +++ b/transformer_engine/pytorch/attention/dot_product_attention/utils.py @@ -21,6 +21,7 @@ import torch.nn.functional as F import transformer_engine_torch as tex import transformer_engine as te +from transformer_engine import te_device_type from transformer_engine.pytorch.cpp_extensions.fused_attn import ( QKVLayout, AttnBiasType, @@ -1193,13 +1194,13 @@ def get_padding_mask( ], dim=0, ) - attention_mask_q = attention_mask_q.to(device="cuda") + attention_mask_q = attention_mask_q.to(device=te_device_type()) if attention_type == "self": attention_mask = attention_mask_q else: attention_mask = ( attention_mask_q, - attention_mask_kv.to(device="cuda"), + attention_mask_kv.to(device=te_device_type()), ) return attention_mask @@ -1318,9 +1319,11 @@ def get_full_mask( actual_seqlens_kv = m[:, 0, 0, :].sum(dim=1) # apply SWA mask - mask = torch.arange(max_seqlen_q, dtype=torch.int32, device="cuda").view( + mask = torch.arange(max_seqlen_q, dtype=torch.int32, device=te_device_type()).view( 1, 1, max_seqlen_q, 1 - ) - torch.arange(max_seqlen_kv, dtype=torch.int32, device="cuda").view(1, 1, 1, max_seqlen_kv) + ) - torch.arange(max_seqlen_kv, dtype=torch.int32, device=te_device_type()).view( + 1, 1, 1, max_seqlen_kv + ) swa_left = None swa_right = None if attn_mask_type == "causal_bottom_right" or ( @@ -1416,7 +1419,7 @@ def get_alibi( m_hat = torch.pow(m_hat_0, torch.arange(1, 1 + 2 * (num_heads - n), 2)) m = torch.cat([m, m_hat]) - _alibi_cache["_alibi_slopes"] = m.to(dtype=torch.float32, device="cuda") + _alibi_cache["_alibi_slopes"] = m.to(dtype=torch.float32, device=te_device_type()) _alibi_cache["_num_heads"] = num_heads _alibi_cache["_alibi_slopes_require_update"] = False @@ -1429,9 +1432,9 @@ def get_alibi( else: raise ValueError("ALiBi slopes cannot exceed 2 dimensions.") - bias = torch.arange(max_seqlen_q, dtype=torch.int32, device="cuda").view( + bias = torch.arange(max_seqlen_q, dtype=torch.int32, device=te_device_type()).view( 1, 1, max_seqlen_q, 1 - ) - torch.arange(max_seqlen_kv, dtype=torch.int32, device="cuda").view( + ) - torch.arange(max_seqlen_kv, dtype=torch.int32, device=te_device_type()).view( 1, 1, 1, max_seqlen_kv ) if actual_seqlens_q is None and actual_seqlens_kv is None: @@ -1451,7 +1454,9 @@ def get_alibi( _alibi_cache["_max_seqlen_q"], _alibi_cache["_max_seqlen_kv"] = max_seqlen_q, max_seqlen_kv _alibi_cache["_bottom_right_alignment"] = bottom_right_alignment bias_dtype = torch.float32 if bias_dtype is None else bias_dtype - _alibi_cache["_alibi_bias"] = bias.contiguous().to(dtype=bias_dtype, device="cuda") + _alibi_cache["_alibi_bias"] = bias.contiguous().to( + dtype=bias_dtype, device=te_device_type() + ) _alibi_cache["_alibi_bias_require_update"] = False return _alibi_cache["_alibi_slopes"], _alibi_cache["_alibi_bias"] @@ -1466,7 +1471,7 @@ def get_cu_seqlens(mask: torch.Tensor) -> torch.Tensor: mask = mask.squeeze(1).squeeze(1) reduced_mask = mask.logical_not().sum(dim=1) cu_seqlens = reduced_mask.cumsum(dim=0).to(torch.int32) - zero = torch.zeros(1, dtype=torch.int32, device="cuda") + zero = torch.zeros(1, dtype=torch.int32, device=te_device_type()) cu_seqlens = torch.cat((zero, cu_seqlens)) return cu_seqlens @@ -1484,7 +1489,7 @@ def get_cu_seqlens_and_indices(mask: torch.Tensor) -> Tuple[torch.Tensor, torch. reduced_mask = mask.logical_not().sum(dim=1) cu_seqlens = reduced_mask.cumsum(dim=0).to(torch.int32) - zero = torch.zeros(1, dtype=torch.int32, device="cuda") + zero = torch.zeros(1, dtype=torch.int32, device=te_device_type()) cu_seqlens = torch.cat((zero, cu_seqlens)) mask = mask.reshape(-1) @@ -1509,7 +1514,12 @@ def get_indices(max_seqlen: int, cu_seqlens: torch.Tensor) -> torch.Tensor: bs = len(cu_seqlens) - 1 seqlens = cu_seqlens[1:] - cu_seqlens[:-1] indices = [i * max_seqlen + ii for i, j in enumerate(seqlens) for ii in range(j)] - indices = torch.Tensor(indices).unsqueeze(1).unsqueeze(1).to(dtype=torch.int64, device="cuda") + indices = ( + torch.Tensor(indices) + .unsqueeze(1) + .unsqueeze(1) + .to(dtype=torch.int64, device=te_device_type()) + ) num_nonzeros = indices.shape[0] pad_amount = bs * max_seqlen - num_nonzeros diff --git a/transformer_engine/pytorch/attention/inference.py b/transformer_engine/pytorch/attention/inference.py index f0ef8d0bd5..fabc491835 100644 --- a/transformer_engine/pytorch/attention/inference.py +++ b/transformer_engine/pytorch/attention/inference.py @@ -11,6 +11,7 @@ import torch import transformer_engine_torch as tex +from transformer_engine import te_device_type from transformer_engine.pytorch.cpp_extensions.fused_attn import QKVFormat __all__ = ["InferenceParams", "KVCacheManager", "NonPagedKVCacheManager", "PagedKVCacheManager"] @@ -626,7 +627,7 @@ def __init__( self.allocated_pages = defaultdict(list) # page table, [batch_size, max_pages_per_seq] self.page_table = torch.zeros( - self.max_batch_size, self.max_pages_per_seq, dtype=torch.int32, device="cuda" + self.max_batch_size, self.max_pages_per_seq, dtype=torch.int32, device=te_device_type() ) def reset(self): diff --git a/transformer_engine/pytorch/attention/multi_head_attention.py b/transformer_engine/pytorch/attention/multi_head_attention.py index b3bda677bb..54c9beb653 100644 --- a/transformer_engine/pytorch/attention/multi_head_attention.py +++ b/transformer_engine/pytorch/attention/multi_head_attention.py @@ -8,6 +8,7 @@ from typing import Callable, List, Optional, Tuple, Union import torch +from transformer_engine import te_device_type from transformer_engine.debug.pytorch.debug_state import TEDebugState from transformer_engine.pytorch.quantization import FP8GlobalStateManager from transformer_engine.pytorch.tensor.float8_tensor import Float8Tensor @@ -255,7 +256,7 @@ def __init__( ub_bulk_wgrad: bool = False, bias: bool = True, normalization: str = "LayerNorm", - device: Union[torch.device, str] = "cuda", + device: Union[torch.device, str] = te_device_type(), qkv_format: str = "sbhd", name: str = None, qk_norm_type: Optional[str] = None, diff --git a/transformer_engine/pytorch/attention/rope.py b/transformer_engine/pytorch/attention/rope.py index cc23d65a3e..bbd5221381 100644 --- a/transformer_engine/pytorch/attention/rope.py +++ b/transformer_engine/pytorch/attention/rope.py @@ -9,6 +9,7 @@ import torch import transformer_engine_torch as tex +from transformer_engine import te_device_type from transformer_engine.pytorch.cpp_extensions.fused_attn import QKVFormat @@ -76,7 +77,7 @@ def forward(self, max_seq_len: int, offset: int = 0): offset: int, default = 0 Fixed offset for frequencies. """ - with torch.autocast(enabled=False, device_type="cuda"): + with torch.autocast(enabled=False, device_type=te_device_type()): seq = ( torch.arange(max_seq_len, device=self.inv_freq.device, dtype=self.inv_freq.dtype) + offset diff --git a/transformer_engine/pytorch/cpp_extensions/gemm.py b/transformer_engine/pytorch/cpp_extensions/gemm.py index a45fafb68a..68c2c20cca 100644 --- a/transformer_engine/pytorch/cpp_extensions/gemm.py +++ b/transformer_engine/pytorch/cpp_extensions/gemm.py @@ -8,6 +8,9 @@ import os import torch import transformer_engine_torch as tex + +from transformer_engine import te_device_type + from ..constants import TE_DType from ..utils import get_sm_count, _empty_tensor @@ -189,7 +192,8 @@ def general_grouped_gemm( sm_count = get_sm_count() if grad and use_bias: grad_bias = [ - torch.empty(B[i].shape[1], dtype=out[0].dtype, device="cuda") for i in range(num_gemms) + torch.empty(B[i].shape[1], dtype=out[0].dtype, device=te_device_type()) + for i in range(num_gemms) ] else: grad_bias = empty_tensors diff --git a/transformer_engine/pytorch/distributed.py b/transformer_engine/pytorch/distributed.py index 5ed73f6783..904f308d1d 100644 --- a/transformer_engine/pytorch/distributed.py +++ b/transformer_engine/pytorch/distributed.py @@ -29,6 +29,8 @@ import transformer_engine_torch as tex +from transformer_engine import te_device_type + from . import torch_version from .utils import ( is_non_tn_fp8_gemm_supported, @@ -90,7 +92,7 @@ def graph_safe_rng_available() -> bool: def _get_cuda_rng_state( - device: Union[int, str, torch.device] = "cuda", + device: Union[int, str, torch.device] = te_device_type(), clone: bool = False, graph_safe: bool = True, ) -> torch.Tensor: @@ -100,7 +102,7 @@ def _get_cuda_rng_state( if isinstance(device, str): device = torch.device(device) elif isinstance(device, int): - device = torch.device("cuda", device) + device = torch.device(te_device_type(), device) idx = device.index if idx is None: idx = torch.cuda.current_device() @@ -122,11 +124,11 @@ def _set_cuda_rng_state( """Sets the random number generator state of the current GPU.""" if device == -1: - device = torch.device("cuda") + device = torch.device(te_device_type()) elif isinstance(device, str): device = torch.device(device) elif isinstance(device, int): - device = torch.device("cuda", device) + device = torch.device(te_device_type(), device) def cb() -> None: idx = device.index @@ -280,10 +282,10 @@ def _get_active_autocast_contexts(): autocast_cached = torch.is_autocast_cache_enabled() if torch_version() >= (2, 4, 0): - gpu_autocast_enabled = torch.is_autocast_enabled("cuda") - gpu_autocast_dtype = torch.get_autocast_dtype("cuda") + gpu_autocast_enabled = torch.is_autocast_enabled(te_device_type()) + gpu_autocast_dtype = torch.get_autocast_dtype(te_device_type()) gpu_autocast_ctx = torch.amp.autocast( - "cuda", + te_device_type(), enabled=gpu_autocast_enabled, dtype=gpu_autocast_dtype, cache_enabled=autocast_cached, @@ -943,7 +945,7 @@ def _all_gather_fp8( out: Float8TensorStorage if quantizer is not None: dtype = torch.float32 - device = "cuda" + device = te_device_type() if isinstance(inp, Float8Tensor): dtype = inp.dtype device = inp.device diff --git a/transformer_engine/pytorch/jit.py b/transformer_engine/pytorch/jit.py index f0f77621e5..32a8deaf45 100644 --- a/transformer_engine/pytorch/jit.py +++ b/transformer_engine/pytorch/jit.py @@ -3,11 +3,15 @@ # See LICENSE for license information. """NVFuser functions and JIT utilities""" + +# pylint: disable=ungrouped-imports + import os from functools import wraps from typing import Callable, Optional, Tuple import torch +from transformer_engine import te_device_type from . import torch_version from .export import is_in_onnx_export_mode from .utils import gpu_autocast_ctx @@ -277,9 +281,13 @@ def warmup_jit_bias_dropout_add( # Save cuda RNG state to ensure warmup does not affect reproducibility. rng_state = torch.cuda.get_rng_state() - inp = torch.rand((seq_length, micro_batch_size, hidden_size), dtype=dtype, device="cuda") - residual = torch.rand((seq_length, micro_batch_size, hidden_size), dtype=dtype, device="cuda") - bias = torch.rand((hidden_size), dtype=dtype, device="cuda") + inp = torch.rand( + (seq_length, micro_batch_size, hidden_size), dtype=dtype, device=te_device_type() + ) + residual = torch.rand( + (seq_length, micro_batch_size, hidden_size), dtype=dtype, device=te_device_type() + ) + bias = torch.rand((hidden_size), dtype=dtype, device=te_device_type()) dropout_rate = 0.1 # Warmup JIT fusions with the input grad_enable state of both forward # prop and recomputation @@ -314,11 +322,11 @@ def warmup_jit_bias_gelu( # Save cuda RNG state to ensure warmup does not affect reproducibility. rng_state = torch.cuda.get_rng_state() - bias = torch.rand(ffn_hidden_size_per_partition, dtype=dtype, device="cuda") + bias = torch.rand(ffn_hidden_size_per_partition, dtype=dtype, device=te_device_type()) inp = torch.rand( (seq_length * micro_batch_size, ffn_hidden_size_per_partition), dtype=dtype, - device="cuda", + device=te_device_type(), ) # Warmup JIT fusions with the input grad_enable state of both forward # prop and recomputation @@ -352,7 +360,7 @@ def warmup_jit_l2normalization( inp = torch.rand( (seq_length * micro_batch_size, hidden_size), dtype=dtype, - device="cuda", + device=te_device_type(), ) eps = 1e-6 # Warmup JIT fusions with the input grad_enable state of both forward diff --git a/transformer_engine/pytorch/module/base.py b/transformer_engine/pytorch/module/base.py index d16455b5b4..06d0de5072 100644 --- a/transformer_engine/pytorch/module/base.py +++ b/transformer_engine/pytorch/module/base.py @@ -19,8 +19,10 @@ import torch.nn.functional as F import transformer_engine_torch as tex +from transformer_engine import te_device_type, te_platform from transformer_engine.common.recipe import Recipe + from ._common import _ParameterInitMeta, noop_cat from ..quantization import ( MXFP8BlockScalingRecipeState, @@ -87,7 +89,9 @@ def get_workspace() -> torch.Tensor: global _cublas_workspace if _cublas_workspace is None: _cublas_workspace = torch.empty( - get_cublas_workspace_size_bytes(), dtype=torch.uint8, device="cuda" + get_cublas_workspace_size_bytes(), + dtype=torch.uint8, + device=te_device_type(), ) return _cublas_workspace @@ -98,7 +102,9 @@ def get_multi_stream_cublas_workspace() -> List[torch.Tensor]: if not _multi_stream_cublas_workspace: for _ in range(tex.get_num_cublas_streams()): _multi_stream_cublas_workspace.append( - torch.empty(get_cublas_workspace_size_bytes(), dtype=torch.uint8, device="cuda") + torch.empty( + get_cublas_workspace_size_bytes(), dtype=torch.uint8, device=te_device_type() + ) ) return _multi_stream_cublas_workspace @@ -111,7 +117,7 @@ def get_dummy_wgrad(shape: list, dtype: torch.dtype, zero=False) -> torch.Tensor _dummy_wgrads[(shape[0], shape[1], dtype)] = torch.empty( shape, dtype=dtype, - device="cuda", + device=te_device_type(), requires_grad=False, ) if zero: @@ -282,7 +288,9 @@ def initialize_ub( elif _cublas_workspace.numel() != get_cublas_workspace_size_bytes() * _NUM_MAX_UB_STREAMS: # This ensures we don't do `.repeat()` on an already expanded workspace _cublas_workspace = torch.empty( - get_cublas_workspace_size_bytes(), dtype=torch.uint8, device="cuda" + get_cublas_workspace_size_bytes(), + dtype=torch.uint8, + device=te_device_type(), ).repeat(_NUM_MAX_UB_STREAMS) # Default buffer precision: AllGather buffers use fp8 when using fp8 recipe @@ -640,7 +648,7 @@ class TransformerEngineBaseModule(torch.nn.Module, ABC): def __init__(self) -> None: super().__init__() - assert torch.cuda.is_available(), "TransformerEngine needs CUDA." + assert te_platform().is_available(), f"TransformerEngine needs {te_device_type()}." self.name = None self.next_iter_when_debug_should_be_run = 0 self.fp8_initialized = False @@ -917,7 +925,7 @@ def set_extra_state(self, state: torch.Tensor) -> None: elif isinstance(state, io.BytesIO): # Deprecated format with io.BytesIO state.seek(0) - state = torch.load(state, map_location="cuda") + state = torch.load(state, map_location=te_device_type()) else: raise RuntimeError("Unsupported checkpoint format.") @@ -1080,7 +1088,9 @@ def prepare_forward( if self.fp8 and in_fp8_activation_recompute_phase(): FP8GlobalStateManager.get_old_fp8_meta_tensors_for_recompute(self.fp8_meta) else: - assert inp.is_cuda, "TransformerEngine needs CUDA." + assert ( + inp.device.type == te_device_type() + ), f"TransformerEngine needs {te_device_type()}." if self.tp_size > 1: assert self.tp_group_initialized, "TP group not initialized." @@ -1257,7 +1267,7 @@ def reset_parameters(self, defer_init: Optional[bool] = False) -> None: for name, param in self.named_parameters(recurse=False): # Ensure parameter is on a real device if param.device == torch.device("meta"): - param = torch.empty_like(param, device="cuda") + param = torch.empty_like(param, device=te_device_type()) # Initialize the parameter values on device init_fn = self.param_init_meta[name].init_fn diff --git a/transformer_engine/pytorch/module/grouped_linear.py b/transformer_engine/pytorch/module/grouped_linear.py index a5bf21ee17..9de94f0ec9 100644 --- a/transformer_engine/pytorch/module/grouped_linear.py +++ b/transformer_engine/pytorch/module/grouped_linear.py @@ -11,7 +11,10 @@ import transformer_engine_torch as tex +from transformer_engine import te_device_type from transformer_engine.common.recipe import Recipe + + from .base import ( get_multi_stream_cublas_workspace, TransformerEngineBaseModule, @@ -581,7 +584,7 @@ def __init__( return_bias: bool = False, params_dtype: Optional[torch.dtype] = None, parallel_mode: Optional[str] = None, - device: Union[torch.device, str] = "cuda", + device: Union[torch.device, str] = te_device_type(), ub_overlap_rs: bool = False, ub_overlap_ag: bool = False, ub_name: Optional[str] = None, diff --git a/transformer_engine/pytorch/module/layernorm_mlp.py b/transformer_engine/pytorch/module/layernorm_mlp.py index a2ddb970af..60db65b0e0 100644 --- a/transformer_engine/pytorch/module/layernorm_mlp.py +++ b/transformer_engine/pytorch/module/layernorm_mlp.py @@ -15,9 +15,12 @@ import transformer_engine_torch as tex +from transformer_engine import te_device_type from transformer_engine.common.recipe import Recipe from transformer_engine.pytorch import torch_version from transformer_engine.pytorch.tensor.utils import is_experimental + + from .base import ( fill_userbuffers_buffer_for_all_gather, get_workspace, @@ -1098,7 +1101,7 @@ def fc2_wgrad_gemm( reduce_scatter_out = None if ctx.ub_overlap_rs_dgrad: reduce_scatter_out = torch.empty( - fc1_dgrad_shape, dtype=ctx.activation_dtype, device="cuda" + fc1_dgrad_shape, dtype=ctx.activation_dtype, device=te_device_type() ) if ctx.ub_bulk_wgrad: gemm_out = ub_obj_fc1_wgrad.get_buffer(local_chunk=False) @@ -1181,7 +1184,7 @@ def fc2_wgrad_gemm( reduce_scatter_out = None if ctx.ub_bulk_wgrad and ub_obj_fc1_wgrad.is_fp8_ubuf(): reduce_scatter_out = torch.empty( - fc1_dgrad_shape, dtype=ctx.activation_dtype, device="cuda" + fc1_dgrad_shape, dtype=ctx.activation_dtype, device=te_device_type() ) # Arguments to include in wgrad GEMM closure diff --git a/transformer_engine/pytorch/ops/basic/activation.py b/transformer_engine/pytorch/ops/basic/activation.py index 8a754c6382..5aa0bc03c0 100644 --- a/transformer_engine/pytorch/ops/basic/activation.py +++ b/transformer_engine/pytorch/ops/basic/activation.py @@ -11,12 +11,16 @@ import torch import transformer_engine_torch as tex + +from transformer_engine import te_device_type + from ...cpu_offload import is_cpu_offload_enabled, mark_activation_offload from ...tensor.float8_tensor import Float8CurrentScalingQuantizer, Quantizer from ...utils import clear_tensor_data from ..op import BasicOperation, OperationContext from .._common import maybe_dequantize + __all__ = [ "GELU", "GEGLU", @@ -92,7 +96,7 @@ def op_forward( # Compute dtype dtype: torch.dtype if torch.is_autocast_enabled(): - dtype = torch.get_autocast_dtype("cuda") + dtype = torch.get_autocast_dtype(te_device_type()) else: dtype = input_.dtype if dtype not in (torch.float32, torch.float16, torch.bfloat16): diff --git a/transformer_engine/pytorch/ops/basic/basic_linear.py b/transformer_engine/pytorch/ops/basic/basic_linear.py index 432d8c134b..18951a316e 100644 --- a/transformer_engine/pytorch/ops/basic/basic_linear.py +++ b/transformer_engine/pytorch/ops/basic/basic_linear.py @@ -12,6 +12,8 @@ import torch +from transformer_engine import te_device_type + from ...cpp_extensions import general_gemm from ...cpu_offload import is_cpu_offload_enabled, mark_activation_offload from ...distributed import ( @@ -967,7 +969,7 @@ def op_forward( # Get autocast dtype if needed if torch.is_autocast_enabled(): - dtype = torch.get_autocast_dtype("cuda") + dtype = torch.get_autocast_dtype(te_device_type()) else: dtype = self.weight.dtype diff --git a/transformer_engine/pytorch/ops/basic/bias.py b/transformer_engine/pytorch/ops/basic/bias.py index 5ec0d2ce5e..e773c35197 100644 --- a/transformer_engine/pytorch/ops/basic/bias.py +++ b/transformer_engine/pytorch/ops/basic/bias.py @@ -10,6 +10,9 @@ import torch import transformer_engine_torch as tex + +from transformer_engine import te_device_type + from ..op import BasicOperation, OperationContext from ...utils import canonicalize_device, canonicalize_dtype from ...tensor import Quantizer @@ -94,7 +97,7 @@ def reset_parameters(self) -> None: # Make sure parameter is initialized bias = self.bias - if bias.device.type != "cuda": + if bias.device.type != te_device_type(): bias = torch.empty_like(bias, device=self.device) else: bias = bias.to(device=self.device) diff --git a/transformer_engine/pytorch/ops/fused/forward_linear_bias_activation.py b/transformer_engine/pytorch/ops/fused/forward_linear_bias_activation.py index 74bd3d1b32..90a16b1d9d 100644 --- a/transformer_engine/pytorch/ops/fused/forward_linear_bias_activation.py +++ b/transformer_engine/pytorch/ops/fused/forward_linear_bias_activation.py @@ -10,6 +10,8 @@ import torch +from transformer_engine import te_device_type + from ...cpu_offload import is_cpu_offload_enabled, mark_activation_offload from ...quantization import FP8GlobalStateManager from ...tensor import Quantizer @@ -95,7 +97,7 @@ def fuser_forward( # Get autocast dtype if needed if torch.is_autocast_enabled(): - dtype = torch.get_autocast_dtype("cuda") + dtype = torch.get_autocast_dtype(te_device_type()) else: dtype = linear_op.weight.dtype diff --git a/transformer_engine/pytorch/ops/fused/forward_linear_bias_add.py b/transformer_engine/pytorch/ops/fused/forward_linear_bias_add.py index 6d5d553391..ab6c2a61b5 100644 --- a/transformer_engine/pytorch/ops/fused/forward_linear_bias_add.py +++ b/transformer_engine/pytorch/ops/fused/forward_linear_bias_add.py @@ -10,6 +10,10 @@ import torch + +from transformer_engine import te_device_type + + from ...cpu_offload import is_cpu_offload_enabled, mark_activation_offload from ...quantization import FP8GlobalStateManager from ...tensor import Quantizer @@ -89,7 +93,7 @@ def fuser_forward( # Get autocast dtype if needed if torch.is_autocast_enabled(): - dtype = torch.get_autocast_dtype("cuda") + dtype = torch.get_autocast_dtype(te_device_type()) else: dtype = linear_op.weight.dtype diff --git a/transformer_engine/pytorch/ops/fused/forward_linear_scale_add.py b/transformer_engine/pytorch/ops/fused/forward_linear_scale_add.py index 24788bcdfb..bfcc1c3f3c 100644 --- a/transformer_engine/pytorch/ops/fused/forward_linear_scale_add.py +++ b/transformer_engine/pytorch/ops/fused/forward_linear_scale_add.py @@ -10,6 +10,8 @@ import torch +from transformer_engine import te_device_type + from ...cpu_offload import is_cpu_offload_enabled, mark_activation_offload from ...quantization import FP8GlobalStateManager from ...tensor import Quantizer @@ -71,7 +73,7 @@ def fuser_forward( # Get autocast dtype if needed if torch.is_autocast_enabled(): - dtype = torch.get_autocast_dtype("cuda") + dtype = torch.get_autocast_dtype(te_device_type()) else: dtype = linear_op.weight.dtype diff --git a/transformer_engine/pytorch/ops/fused/userbuffers_backward_linear.py b/transformer_engine/pytorch/ops/fused/userbuffers_backward_linear.py index d95b2298fe..0759abbc0c 100644 --- a/transformer_engine/pytorch/ops/fused/userbuffers_backward_linear.py +++ b/transformer_engine/pytorch/ops/fused/userbuffers_backward_linear.py @@ -11,6 +11,9 @@ import torch from transformer_engine_torch import CommOverlapType, bulk_overlap_ag_with_external_gemm + +from transformer_engine import te_device_type + from ...cpp_extensions import general_gemm from ...distributed import get_distributed_world_size from ...module.base import ( @@ -176,7 +179,7 @@ def _functional_backward( else: device = grad_output.device device = canonicalize_device(device) - if device.type != "cuda": + if device.type != te_device_type(): raise ValueError(f"Only CUDA devices are supported (got {device})") # Check datatype diff --git a/transformer_engine/pytorch/ops/fused/userbuffers_forward_linear.py b/transformer_engine/pytorch/ops/fused/userbuffers_forward_linear.py index e20de53da3..08e7d92d42 100644 --- a/transformer_engine/pytorch/ops/fused/userbuffers_forward_linear.py +++ b/transformer_engine/pytorch/ops/fused/userbuffers_forward_linear.py @@ -11,6 +11,7 @@ import torch from transformer_engine_torch import CommOverlapType +from transformer_engine import te_device_type from ...cpp_extensions import general_gemm from ...cpu_offload import is_cpu_offload_enabled, mark_activation_offload from ...distributed import get_distributed_world_size @@ -156,7 +157,7 @@ def _functional_forward( """ # Check device - if device.type != "cuda": + if device.type != te_device_type(): raise ValueError(f"Only CUDA devices are supported (got {device})") # Check datatype @@ -322,7 +323,7 @@ def fuser_forward( # Get autocast dtype if needed if torch.is_autocast_enabled(): - dtype = torch.get_autocast_dtype("cuda") + dtype = torch.get_autocast_dtype(te_device_type()) else: dtype = linear_op.weight.dtype diff --git a/transformer_engine/pytorch/optimizers/fused_adam.py b/transformer_engine/pytorch/optimizers/fused_adam.py index 18f7e2031a..3d71f4ff63 100644 --- a/transformer_engine/pytorch/optimizers/fused_adam.py +++ b/transformer_engine/pytorch/optimizers/fused_adam.py @@ -12,6 +12,7 @@ import torch import transformer_engine_torch as tex +from transformer_engine import te_device_type from transformer_engine.pytorch.tensor.float8_tensor import Float8Tensor, Float8Quantizer from .multi_tensor_apply import multi_tensor_applier @@ -178,7 +179,7 @@ def __init__( self._step_supports_amp_scaling = True # Skip buffer - self._dummy_overflow_buf = torch.tensor([0], dtype=torch.int, device="cuda") + self._dummy_overflow_buf = torch.tensor([0], dtype=torch.int, device=te_device_type()) self.multi_tensor_adam = tex.multi_tensor_adam self.multi_tensor_adam_param_remainder = tex.multi_tensor_adam_param_remainder self.multi_tensor_adam_fp8 = tex.multi_tensor_adam_fp8 diff --git a/transformer_engine/pytorch/permutation.py b/transformer_engine/pytorch/permutation.py index ea3e67a57c..23dbbf3598 100644 --- a/transformer_engine/pytorch/permutation.py +++ b/transformer_engine/pytorch/permutation.py @@ -8,6 +8,7 @@ import torch import transformer_engine_torch as tex +from transformer_engine import te_device_type import transformer_engine.pytorch.triton.permutation as triton_permutation from transformer_engine.pytorch.constants import TE_DType from transformer_engine.pytorch.tensor.quantized_tensor import QuantizedTensor @@ -15,6 +16,7 @@ from transformer_engine.pytorch.tensor.float8_blockwise_tensor import Float8BlockwiseQTensor from transformer_engine.pytorch.tensor.mxfp8_tensor import MXFP8Tensor + __all__ = [ "moe_permute", "moe_unpermute", @@ -42,8 +44,8 @@ def forward( return inp, torch.tensor([], device=inp.device) # Device check - assert inp.is_cuda, "TransformerEngine needs CUDA." - assert index.is_cuda, "TransformerEngine needs CUDA." + assert inp.device.type == te_device_type(), f"TransformerEngine needs {te_device_type()}." + assert index.device.type == te_device_type(), f"TransformerEngine needs {te_device_type()}." # Shape check assert inp.size(0) == index.size(0), "Permute not possible" @@ -119,7 +121,9 @@ def forward( # None probs check if probs is not None: - assert probs.is_cuda, "TransformerEngine needs CUDA." + assert ( + probs.device.type == te_device_type() + ), f"TransformerEngine needs {te_device_type()}." if probs.dtype != torch.float32: warnings.warn( @@ -136,8 +140,10 @@ def forward( probs = torch.empty(0) # Device check - assert inp.is_cuda, "TransformerEngine needs CUDA." - assert row_id_map.is_cuda, "TransformerEngine needs CUDA." + assert inp.device.type == te_device_type(), f"TransformerEngine needs {te_device_type()}." + assert ( + row_id_map.device.type == te_device_type() + ), f"TransformerEngine needs {te_device_type()}." # Data type check dtype = TE_DType[inp.dtype] @@ -197,10 +203,14 @@ def forward( ctx.probs = probs return inp, torch.tensor([], device=inp.device), torch.tensor([], device=inp.device) - assert inp.is_cuda, "TransformerEngine needs CUDA." - assert routing_map.is_cuda, "TransformerEngine needs CUDA." + assert inp.device.type == te_device_type(), f"TransformerEngine needs {te_device_type()}." + assert ( + routing_map.device.type == te_device_type() + ), f"TransformerEngine needs {te_device_type()}." if probs is not None: - assert probs.is_cuda, "TransformerEngine needs CUDA." + assert ( + probs.device.type == te_device_type() + ), f"TransformerEngine needs {te_device_type()}." assert inp.size(0) == routing_map.size(0), "Permute not possible" num_tokens, hidden_size = inp.size() @@ -353,11 +363,15 @@ def forward( with_probs = merging_probs is not None if with_probs: - assert merging_probs.is_cuda, "TransformerEngine needs CUDA." + assert ( + merging_probs.device.type == te_device_type() + ), f"TransformerEngine needs {te_device_type()}." # Device check - assert inp.is_cuda, "TransformerEngine needs CUDA." - assert row_id_map.is_cuda, "TransformerEngine needs CUDA." + assert inp.device.type == te_device_type(), f"TransformerEngine needs {te_device_type()}." + assert ( + row_id_map.device.type == te_device_type() + ), f"TransformerEngine needs {te_device_type()}." assert not isinstance( inp, QuantizedTensor @@ -635,11 +649,17 @@ def forward( if not inp.numel(): return inp, probs - assert inp.is_cuda, "TransformerEngine needs CUDA." - assert split_sizes.is_cuda, "TransformerEngine needs CUDA." - assert sorted_idxs.is_cuda, "TransformerEngine needs CUDA." + assert inp.device.type == te_device_type(), f"TransformerEngine needs {te_device_type()}." + assert ( + split_sizes.device.type == te_device_type() + ), f"TransformerEngine needs {te_device_type()}." + assert ( + sorted_idxs.device.type == te_device_type() + ), f"TransformerEngine needs {te_device_type()}." if probs is not None: - assert probs.is_cuda, "TransformerEngine needs CUDA." + assert ( + probs.device.type == te_device_type() + ), f"TransformerEngine needs {te_device_type()}." num_tokens, hidden_size = inp.shape num_splits = split_sizes.size(0) diff --git a/transformer_engine/pytorch/quantization.py b/transformer_engine/pytorch/quantization.py index 030370b9db..9ea48964ea 100644 --- a/transformer_engine/pytorch/quantization.py +++ b/transformer_engine/pytorch/quantization.py @@ -16,6 +16,7 @@ import torch import transformer_engine_torch as tex +from transformer_engine import te_device_type from transformer_engine.common.recipe import ( Recipe, DelayedScaling, @@ -27,6 +28,7 @@ CustomRecipe, ) + from .constants import dist_group_type from .utils import get_device_compute_capability from .jit import jit_fuser @@ -279,7 +281,9 @@ def reset(cls) -> None: def set_skip_fp8_weight_update_tensor(cls, skip: bool) -> None: """`skip_fp8_weight_update_tensor` inplace setter.""" if cls.skip_fp8_weight_update_tensor is None: - cls.skip_fp8_weight_update_tensor = torch.empty(1, dtype=torch.float32, device="cuda") + cls.skip_fp8_weight_update_tensor = torch.empty( + 1, dtype=torch.float32, device=te_device_type() + ) cls.skip_fp8_weight_update_tensor.fill_(skip) @classmethod @@ -1067,7 +1071,7 @@ def __init__( # Allocate buffers if device is None: - device = torch.device("cuda") + device = torch.device(te_device_type()) self.scale = torch.ones(num_quantizers, dtype=torch.float32, device=device) self.amax_history = torch.zeros( recipe.amax_history_len, @@ -1113,7 +1117,7 @@ def __init__( # Allocate buffers if device is None: - device = torch.device("cuda") + device = torch.device(te_device_type()) self.device = device def make_quantizers(self) -> list: @@ -1153,7 +1157,7 @@ def __init__( # Allocate buffers if device is None: - device = torch.device("cuda") + device = torch.device(te_device_type()) def make_quantizers(self) -> list: # TODO(ksivamani); Find better design for this, adding here to avoid circular import. @@ -1192,7 +1196,7 @@ def __init__( # Allocate buffers if device is None: - device = torch.device("cuda") + device = torch.device(te_device_type()) self.device = device def make_quantizers(self) -> list: @@ -1293,7 +1297,7 @@ def __init__( # Allocate buffers if device is None: - device = torch.device("cuda") + device = torch.device(te_device_type()) def make_quantizers(self) -> list: from .tensor.nvfp4_tensor import NVFP4Quantizer @@ -1363,7 +1367,7 @@ def __init__( self.mode = mode self.num_quantizers = num_quantizers if device is None: - device = torch.device("cuda") + device = torch.device(te_device_type()) self.device = device if getattr(recipe, "qfactory", None) is None: diff --git a/transformer_engine/pytorch/tensor/float8_blockwise_tensor.py b/transformer_engine/pytorch/tensor/float8_blockwise_tensor.py index 48762499b9..c752501848 100644 --- a/transformer_engine/pytorch/tensor/float8_blockwise_tensor.py +++ b/transformer_engine/pytorch/tensor/float8_blockwise_tensor.py @@ -12,6 +12,7 @@ from transformer_engine_torch import DType as TE_DType from transformer_engine_torch import Float8BlockScaleTensorFormat +from transformer_engine import te_device_type from transformer_engine.common.recipe import Float8BlockScaling, Recipe from .storage.float8_blockwise_tensor_storage import Float8BlockwiseQTensorStorage from .quantized_tensor import ( @@ -220,7 +221,7 @@ def make_empty( ) -> Float8BlockwiseQTensor: """Construct quantized tensor with uninitialized data""" if device is None: - device = torch.device("cuda") + device = torch.device(te_device_type()) data_format = ( tex.Float8BlockScaleTensorFormat.COMPACT @@ -451,7 +452,7 @@ def __torch_dispatch__(cls, func, types, args, kwargs=None): qt._rowwise_scale_inv, qt._columnwise_scale_inv, ): - if t is not None and t.is_cuda: + if t is not None and t.device.type == te_device_type(): t.record_stream(stream) return None @@ -542,7 +543,7 @@ def _set_data(self, tensor: torch.Tensor) -> None: """ # Tensor device - new_device = tensor.device if tensor.is_cuda else self.device + new_device = tensor.device if tensor.device.type == te_device_type() else self.device def _set_from_tensor(dst: Float8BlockwiseQTensor, src: Float8BlockwiseQTensor): dst._rowwise_data = src._rowwise_data diff --git a/transformer_engine/pytorch/tensor/float8_tensor.py b/transformer_engine/pytorch/tensor/float8_tensor.py index a4e68e53b0..ea88c7e3f2 100644 --- a/transformer_engine/pytorch/tensor/float8_tensor.py +++ b/transformer_engine/pytorch/tensor/float8_tensor.py @@ -11,6 +11,7 @@ import transformer_engine_torch as tex from transformer_engine_torch import DType as TE_DType +from transformer_engine import te_device_type from transformer_engine.common.recipe import DelayedScaling, Float8CurrentScaling, Recipe from ..utils import canonicalize_process_group, devices_match from .storage.float8_tensor_storage import Float8TensorStorage, _FromFloat8Func @@ -108,7 +109,7 @@ def make_empty( # Canonicalize tensor attributes if device is None: - device = torch.device("cuda") + device = torch.device(te_device_type()) # Allocate FP8 data data = torch.empty(shape, dtype=torch.uint8, device=device) @@ -294,7 +295,7 @@ def make_empty( # Canonicalize tensor attributes if device is None: - device = torch.device("cuda") + device = torch.device(te_device_type()) # Allocate FP8 data data = torch.empty(shape, dtype=torch.uint8, device=device) @@ -682,7 +683,7 @@ def _set_data(self, tensor: torch.Tensor) -> None: """ # Tensor device - new_device = tensor.device if tensor.is_cuda else self.device + new_device = tensor.device if tensor.device.type == te_device_type() else self.device if not devices_match(new_device, tensor.device): tensor = tensor.to(device=new_device) diff --git a/transformer_engine/pytorch/tensor/mxfp8_tensor.py b/transformer_engine/pytorch/tensor/mxfp8_tensor.py index 700de24c4e..c8dda346e9 100644 --- a/transformer_engine/pytorch/tensor/mxfp8_tensor.py +++ b/transformer_engine/pytorch/tensor/mxfp8_tensor.py @@ -12,6 +12,7 @@ import transformer_engine_torch as tex from transformer_engine_torch import DType as TE_DType +from transformer_engine import te_device_type from transformer_engine.common.recipe import MXFP8BlockScaling, Recipe from ..constants import MXFP8_BLOCK_SCALING_SIZE from ..utils import devices_match, round_up_to_nearest_multiple @@ -96,7 +97,7 @@ def make_empty( # Canonicalize tensor attributes if device is None: - device = torch.device("cuda") + device = torch.device(te_device_type()) assert ( shape[-1] % MXFP8_BLOCK_SCALING_SIZE == 0 @@ -403,7 +404,7 @@ def _set_data(self, tensor: torch.Tensor) -> None: """ # Tensor device - new_device = tensor.device if tensor.is_cuda else self.device + new_device = tensor.device if tensor.device.type == te_device_type() else self.device if not devices_match(new_device, tensor.device): tensor = tensor.to(device=new_device) diff --git a/transformer_engine/pytorch/tensor/nvfp4_tensor.py b/transformer_engine/pytorch/tensor/nvfp4_tensor.py index ca2154f554..a62873cf7c 100644 --- a/transformer_engine/pytorch/tensor/nvfp4_tensor.py +++ b/transformer_engine/pytorch/tensor/nvfp4_tensor.py @@ -13,6 +13,7 @@ import transformer_engine_torch as tex from transformer_engine_torch import DType as TE_DType +from transformer_engine import te_device_type from transformer_engine.common.recipe import NVFP4BlockScaling, Recipe from ..constants import NVFP4_BLOCK_SCALING_SIZE, dist_group_type from ..utils import ( @@ -96,7 +97,7 @@ def get_rht_matrix(with_random_sign_mask: bool) -> torch.Tensor: signs = get_no_random_sign_vector() sign_matrix = signs * torch.eye(hadamard_dimension, dtype=torch.float32) rht_matrix = sign_matrix @ get_hadamard_matrix(hadamard_dimension) - return rht_matrix.to(dtype=torch.bfloat16).cuda() + return rht_matrix.to(dtype=torch.bfloat16).to(te_device_type()) @functools.lru_cache(maxsize=None) @@ -267,7 +268,7 @@ def make_empty( # Canonicalize tensor attributes if device is None: - device = torch.device("cuda") + device = torch.device(te_device_type()) assert shape[-1] % NVFP4_BLOCK_SCALING_SIZE == 0, ( f"Incorrect shape {shape} for NVFP4. Tensor dims must be divisible by" @@ -617,7 +618,7 @@ def _set_data(self, tensor: torch.Tensor) -> None: """ # Tensor device - new_device = tensor.device if tensor.is_cuda else self.device + new_device = tensor.device if tensor.device.type == te_device_type() else self.device if not devices_match(new_device, tensor.device): tensor = tensor.to(device=new_device) diff --git a/transformer_engine/pytorch/transformer.py b/transformer_engine/pytorch/transformer.py index 8a032b2f55..b59f7276bd 100644 --- a/transformer_engine/pytorch/transformer.py +++ b/transformer_engine/pytorch/transformer.py @@ -10,6 +10,7 @@ import torch +from transformer_engine import te_device_type from transformer_engine.pytorch import torch_version from transformer_engine.pytorch.module import LayerNormMLP, LayerNorm, RMSNorm from transformer_engine.debug.pytorch.debug_state import TEDebugState @@ -311,7 +312,7 @@ def __init__( bias: bool = True, activation: str = "gelu", normalization: str = "LayerNorm", - device: Union[torch.device, str] = "cuda", + device: Union[torch.device, str] = te_device_type(), attn_input_format: str = "sbhd", name: str = None, qk_norm_type: Optional[str] = None, diff --git a/transformer_engine/pytorch/triton/permutation.py b/transformer_engine/pytorch/triton/permutation.py index 6292acb69b..aa1260aeac 100644 --- a/transformer_engine/pytorch/triton/permutation.py +++ b/transformer_engine/pytorch/triton/permutation.py @@ -13,6 +13,7 @@ from triton.language import core from triton.language.standard import _log2 +from transformer_engine import te_device_type # The following three argsort related kernels are adapted from # the issue https://github.com/triton-lang/triton/issues/3698 @@ -218,10 +219,12 @@ def make_row_id_map( The [num_experts, num_experts + n_routed) items are the indices of the experts corresponding to the first n_routed row indices above. """ - row_id_map = torch.empty((num_tokens, num_experts * 2 + 1), dtype=torch.int32, device="cuda") + row_id_map = torch.empty( + (num_tokens, num_experts * 2 + 1), dtype=torch.int32, device=te_device_type() + ) block_size = 1024 grid = (num_experts, triton.cdiv(num_tokens, block_size)) - workspace_tensor = torch.empty(grid, dtype=torch.int32, device="cuda") + workspace_tensor = torch.empty(grid, dtype=torch.int32, device=te_device_type()) # supposing num_tokens == 5, num_experts == 3, block_size == 3 # and we have a routing_map like this: @@ -419,15 +422,15 @@ def permute_with_mask_map( scale_hidden_dim: int Hidden size of the scale tensor. """ - output = torch.empty((num_out_tokens, hidden_size), dtype=inp.dtype, device="cuda") + output = torch.empty((num_out_tokens, hidden_size), dtype=inp.dtype, device=te_device_type()) if probs is not None: - permuted_probs = torch.empty((num_out_tokens,), dtype=probs.dtype, device="cuda") + permuted_probs = torch.empty((num_out_tokens,), dtype=probs.dtype, device=te_device_type()) else: permuted_probs = None if scale is not None: permuted_scale = torch.empty( - (num_out_tokens, scale_hidden_dim), dtype=scale.dtype, device="cuda" + (num_out_tokens, scale_hidden_dim), dtype=scale.dtype, device=te_device_type() ) else: permuted_scale = None @@ -603,10 +606,10 @@ def unpermute_with_mask_map( hidden_size: int Hidden size of the permuted tensor. """ - output = torch.empty((num_tokens, hidden_size), dtype=inp.dtype, device="cuda") + output = torch.empty((num_tokens, hidden_size), dtype=inp.dtype, device=te_device_type()) if permuted_probs is not None: unpermuted_probs = torch.empty( - (num_tokens, num_experts), dtype=permuted_probs.dtype, device="cuda" + (num_tokens, num_experts), dtype=permuted_probs.dtype, device=te_device_type() ) else: unpermuted_probs = None @@ -776,10 +779,10 @@ def unpermute_with_mask_map_bwd_with_merging_probs( Hidden size of the output tensor. """ act_grad = torch.empty( - (num_out_tokens, hidden_size), dtype=fwd_output_grad.dtype, device="cuda" + (num_out_tokens, hidden_size), dtype=fwd_output_grad.dtype, device=te_device_type() ) merging_probs_grad = torch.empty( - (num_tokens, num_experts), dtype=merging_probs.dtype, device="cuda" + (num_tokens, num_experts), dtype=merging_probs.dtype, device=te_device_type() ) grid = (num_tokens,) _unpermute_bwd_with_merging_probs_kernel[grid]( @@ -869,7 +872,7 @@ def make_chunk_sort_map( num_splits: int Number of splits of split_sizes and sorted_indices. """ - row_id_map = torch.empty((num_tokens,), dtype=torch.int32, device="cuda") + row_id_map = torch.empty((num_tokens,), dtype=torch.int32, device=te_device_type()) grid = (num_tokens,) _make_chunk_sort_map_kernel[grid]( split_sizes, @@ -968,9 +971,9 @@ def sort_chunks_by_map( is_forward: bool Whether the sort is for forward or backward. """ - output = torch.empty((num_tokens, hidden_size), dtype=inp.dtype, device="cuda") + output = torch.empty((num_tokens, hidden_size), dtype=inp.dtype, device=te_device_type()) if probs is not None: - permuted_probs = torch.empty((num_tokens,), dtype=probs.dtype, device="cuda") + permuted_probs = torch.empty((num_tokens,), dtype=probs.dtype, device=te_device_type()) else: permuted_probs = None # pylint: disable=unnecessary-lambda-assignment diff --git a/transformer_engine/pytorch/utils.py b/transformer_engine/pytorch/utils.py index 2be0aed4a8..7d237ac3da 100644 --- a/transformer_engine/pytorch/utils.py +++ b/transformer_engine/pytorch/utils.py @@ -11,6 +11,8 @@ import numpy as np import torch +from transformer_engine import te_device_type + from . import torch_version from .tensor.quantized_tensor import Quantizer from ..debug.pytorch.debug_quantization import DebugQuantizedTensor @@ -30,7 +32,8 @@ def requires_grad(*tensors: Tuple[Optional[torch.Tensor], ...]) -> None: @functools.lru_cache(maxsize=None) def _empty_tensor() -> torch.Tensor: """Get tensor with no entries and no data""" - return torch.Tensor().cuda() + + return torch.Tensor().to(device=te_device_type()) def clear_tensor_data(*tensors: Tuple[Optional[torch.Tensor], ...]) -> None: @@ -516,12 +519,12 @@ def canonicalize_device(device: Optional[torch.device | str]) -> torch.device: if device is None: # Use default CUDA device device = torch.get_default_device() - if device.type != "cuda": - device = torch.device("cuda", torch.cuda.current_device()) + if device.type != te_device_type(): + device = torch.device(te_device_type(), torch.cuda.current_device()) elif not isinstance(device, torch.device): device = torch.device(device) - if device.type == "cuda" and device.index is None: - device = torch.device("cuda", torch.cuda.current_device()) + if device.type == te_device_type() and device.index is None: + device = torch.device(te_device_type(), torch.cuda.current_device()) return device @@ -543,7 +546,7 @@ def devices_match(device1: torch.device, device2: torch.device) -> bool: device2 = torch.device(device2) if device1.type != device2.type: return False - if device1.type == "cuda": + if device1.type == te_device_type(): index1 = device1.index index2 = device2.index if index1 == index2: @@ -657,12 +660,12 @@ def canonicalize_process_group( def torch_get_autocast_gpu_dtype() -> torch.dtype: """Get PyTorch autocast GPU dtype.""" if torch_version() >= (2, 4, 0): - return torch.get_autocast_dtype("cuda") + return torch.get_autocast_dtype(te_device_type()) return torch.get_autocast_gpu_dtype() if torch_version() >= (2, 4, 0): - gpu_autocast_ctx = functools.partial(torch.amp.autocast, device_type="cuda") + gpu_autocast_ctx = functools.partial(torch.amp.autocast, device_type=te_device_type()) else: gpu_autocast_ctx = torch.cuda.amp.autocast @@ -759,7 +762,7 @@ def convert_to_torch_tensor(tensor: Union[_WeakRefTensor, torch.Tensor]) -> torc if isinstance(x, torch.Tensor): return ( convert_to_torch_tensor(_WeakRefTensor(x.data_ptr(), x.dtype, x.shape)) - if x.is_cuda + if x.device.type == te_device_type() else x ) if isinstance(x, tuple):