leavelet/cutlass
1
0
Fork 0
Code Issues Pull Requests Actions 30 Packages Projects Releases Wiki Activity

v4.3.5 update. (#2934)

Browse Source 
* v4.3.5 update.

* Update copyright to 2026
This commit is contained in:
Junkai-Wu
2026-01-08 15:02:56 -05:00
committed by GitHub
parent f86feb0aa8
commit 0d2b201e8c
2788 changed files with 3172 additions and 3071 deletions
Download Patch File Download Diff File Expand all files Collapse all files
python/CuTeDSL/cutlass/__init__.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/__init__.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/_mlir_helpers/__init__.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/_mlir_helpers/arith.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/_mlir_helpers/gpu.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/_mlir_helpers/lru_cache_ir.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/_mlir_helpers/op.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/arch.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/ast_helpers.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/ast_preprocessor.py
+257 -136
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
@@ -40,7 +40,7 @@ import os
import sys
import textwrap
import warnings
from dataclasses import dataclass
from dataclasses import dataclass, field
from typing import List, Set, Dict, Any, Callable, Optional
from types import ModuleType
from collections import OrderedDict
@@ -161,6 +161,118 @@ class ScopeManager:
self.scopes.pop()
@dataclass
class SessionData:
"""
Session data for the DSL preprocessor.
"""
counter: int = 0 # Unique function names for multiple loops
scope_manager: ScopeManager = field(default_factory=ScopeManager.create)
function_counter: int = 0
function_name: str = "<unknown function>"
class_name: Optional[str] = None
file_name: str = "<unknown filename>"
function_depth: int = 0
local_closures: set[str] = field(default_factory=set)
function_globals: Optional[dict[str, Any]] = None
import_top_module: bool = False
def _create_module_attribute(
func_name,
*,
use_base_dsl=True,
submodule_name="ast_helpers",
lineno=None,
col_offset=None,
):
"""Creates an AST node representing a qualified attribute access to a function in a module or submodule.
:param func_name: The attribute or function name to access
:type func_name: str
:param top_module_name: The top-level module name, defaults to "_dsl_"
:type top_module_name: str, optional
:param submodule_name: The submodule name to access within the top module,
defaults to "ast_helpers"
:type submodule_name: str, optional
:param lineno: The line number to use for AST node location, defaults to None
:type lineno: int, optional
:param col_offset: The column offset to use for AST node location, defaults to None
:type col_offset: int, optional
:return: An AST Attribute node corresponding to the desired attribute access,
with optional location info
:rtype: ast.Attribute
"""
# If we simply copy location from origin node, it contains a way to wide range, which cause location in traceback to be wrong.
def set_location(node, lineno, col_offset):
if lineno is None or col_offset is None:
return
node.lineno = lineno
node.end_lineno = lineno
node.col_offset = col_offset
node.end_col_offset = col_offset
base = ast.Name(
id="__base_dsl__" if use_base_dsl else "__module_dsl__", ctx=ast.Load()
)
set_location(base, lineno, col_offset)
if submodule_name:
base = ast.Attribute(value=base, attr=submodule_name, ctx=ast.Load())
set_location(base, lineno, col_offset)
node = ast.Attribute(value=base, attr=func_name, ctx=ast.Load())
set_location(node, lineno, col_offset)
return node
class DSLPreprocessorSession:
"""Context manager for managing a DSL preprocessor session.
This context manager is used to ensure that each preprocessing operation
(typically a transformation of a Python AST via the DSL preprocessor)
is performed within a well-defined session. When entering the context,
it initializes session-specific resources or state by calling
`_start_session()` on the provided DSL object. Upon exit, it performs
appropriate cleanup by calling `_end_session()`.
Example usage::
with DSLPreprocessorSession(dsl_object):
# perform AST transformations or other preprocessing actions
:param dsl_object: An instance of a DSL object that implements
`_start_session()` and `_end_session()` methods to manage the
session state
:type dsl_object: DSLPreprocessor
"""
def __init__(self, dsl_object):
self.dsl_object = dsl_object
def __enter__(self):
"""Starts the DSL preprocessor session.
:return: The DSL object for use within the context
:rtype: DSLPreprocessor
"""
self.dsl_object._start_session()
# Let `with preprocessor.get_session() as p:` keep using `p` as the preprocessor.
return self.dsl_object
def __exit__(self, exc_type, exc_value, traceback):
"""Ends the DSL preprocessor session.
:param exc_type: The exception type if an exception was raised in the context
:type exc_type: type, optional
:param exc_value: The exception value if an exception was raised in the context
:type exc_value: Exception, optional
:param traceback: The traceback if an exception was raised in the context
:type traceback: traceback, optional
"""
self.dsl_object._end_session()
class DSLPreprocessor(ast.NodeTransformer):
"""
A preprocessor for transforming Python ASTs. It supports:
@@ -185,44 +297,42 @@ class DSLPreprocessor(ast.NodeTransformer):
def __init__(self, client_module_name):
super().__init__()
self.counter = 0 # Unique function names for multiple loops
self.scope_manager = ScopeManager.create()
# Persistent state
self.processed_functions = set()
self.function_counter = 0
self.function_name = "<unknown function>"
self.class_name = None
self.file_name = "<unknown filename>"
self.function_depth = 0
self.local_closures = set()
self.function_globals = None
self.client_module_name = client_module_name
self.import_top_module = False
self.module_cache = {}
self._session_data = None
def _create_module_attribute(
self,
func_name,
*,
top_module_name="_dsl_",
submodule_name="ast_helpers",
lineno=None,
col_offset=None,
):
# If we simply copy location from origin node, it contains a way to wide range, which cause location in traceback to be wrong.
def set_location(node, lineno, col_offset):
if lineno and col_offset:
node.lineno = lineno
node.end_lineno = lineno
node.col_offset = col_offset
node.end_col_offset = col_offset
def _start_session(self):
"""
Starts a new preprocessing session by initializing session data.
base = ast.Name(id=top_module_name, ctx=ast.Load())
set_location(base, lineno, col_offset)
if submodule_name:
base = ast.Attribute(value=base, attr=submodule_name, ctx=ast.Load())
set_location(base, lineno, col_offset)
node = ast.Attribute(value=base, attr=func_name, ctx=ast.Load())
set_location(node, lineno, col_offset)
return node
This method sets up a fresh SessionData instance for use during
AST transformations. It must be called before performing any
preprocessing actions that require access to context-specific
information during the transformation of a function's AST.
"""
self._session_data = SessionData()
def _end_session(self):
"""
Ends the current preprocessing session and clears session data.
This method resets the session-specific data, marking the end of
a preprocessing context. It should be called after all necessary
AST processing is complete to ensure no stale context remains.
"""
self._session_data = None
def get_session(self):
return DSLPreprocessorSession(dsl_object=self)
@property
def session_data(self):
assert self._session_data is not None, (
"Please start a session before accessing session data"
)
return self._session_data
def _get_imports_from_ast(self, node, module):
"""
@@ -328,7 +438,7 @@ class DSLPreprocessor(ast.NodeTransformer):
@contextlib.contextmanager
def local_import(module_path):
# Directory where some local import might happen:
local_dir = os.path.dirname(self.file_name)
local_dir = os.path.dirname(self.session_data.file_name)
# Momentarily insert the directory where the local import
# used to happen, so the import can find the module.
sys.path.insert(0, local_dir)
@@ -389,6 +499,7 @@ class DSLPreprocessor(ast.NodeTransformer):
self.exec_imports(import_info.finally_imports, exec_globals)
def exec(self, function_name, original_function, code_object, exec_globals):
"""Requires an active DSL preprocessor session."""
# Get imports from the original module
module_imports = self._get_module_imports(original_function)
@@ -463,11 +574,17 @@ class DSLPreprocessor(ast.NodeTransformer):
# Step 3. Import cutlass and base_dsl
top_module_name = ".".join(self.client_module_name)
import_stmts = []
if self.import_top_module:
import_stmts.append(ast.Import(names=[ast.alias(name=top_module_name)]))
if self.session_data.import_top_module:
import_stmts.append(
ast.Import(
names=[ast.alias(name=top_module_name, asname="__module_dsl__")]
)
)
import_stmts.append(
ast.Import(
names=[ast.alias(name=f"{top_module_name}.base_dsl", asname="_dsl_")]
names=[
ast.alias(name=f"{top_module_name}.base_dsl", asname="__base_dsl__")
]
)
)
@@ -584,9 +701,13 @@ class DSLPreprocessor(ast.NodeTransformer):
if not checker.has_early_exit:
return
raise DSLAstPreprocessorError(
message=f"Early exit ({checker.early_exit_type}) is not allowed in `{self.function_name}`"
+ (f" in `{self.class_name}`" if self.class_name else ""),
filename=self.file_name,
message=f"Early exit ({checker.early_exit_type}) is not allowed in `{self.session_data.function_name}`"
+ (
f" in `{self.session_data.class_name}`"
if self.session_data.class_name
else ""
),
filename=self.session_data.file_name,
snippet=ast.unparse(tree),
suggestion=(
"If predicates are constant expression, write like "
@@ -633,13 +754,14 @@ class DSLPreprocessor(ast.NodeTransformer):
def transform(self, original_function, exec_globals):
"""
Transforms the provided function using the preprocessor.
Requires an active DSL preprocessor session.
"""
self.file_name = inspect.getsourcefile(original_function)
self.function_globals = exec_globals
self.session_data.file_name = inspect.getsourcefile(original_function)
self.session_data.function_globals = exec_globals
transformed_tree = self.transform_function(
original_function.__name__, original_function
)
self.function_globals = None
self.session_data.function_globals = None
unified_tree = ast.Module(body=transformed_tree, type_ignores=[])
unified_tree = ast.fix_missing_locations(unified_tree)
@@ -656,9 +778,7 @@ class DSLPreprocessor(ast.NodeTransformer):
# we need orderedset to keep the insertion order the same. otherwise generated IR is different each time
write_args = OrderedSet()
invoked_args = OrderedSet()
local_closure = self.local_closures
file_name = self.file_name
region_node = node
local_closure = self.session_data.local_closures
called_closures = OrderedSet()
class RegionAnalyzer(ast.NodeVisitor):
@@ -767,7 +887,8 @@ class DSLPreprocessor(ast.NodeTransformer):
return self.visit(args[0]), self.visit(args[1]), self.visit(args[2]), True
else:
raise DSLAstPreprocessorError(
"Unsupported number of arguments in range", filename=self.file_name
"Unsupported number of arguments in range",
filename=self.session_data.file_name,
)
def extract_unroll_args(self, iter_node):
@@ -790,7 +911,7 @@ class DSLPreprocessor(ast.NodeTransformer):
self.issue_deprecation_warning(
message="pipelining is deprecated, use prefetch_stages instead",
category=DeprecationWarning,
filename=self.file_name,
filename=self.session_data.file_name,
lineno=iter_node.lineno,
)
return keywords.get("pipelining", ast.Constant(value=None))
@@ -841,7 +962,7 @@ class DSLPreprocessor(ast.NodeTransformer):
# Define the decorator with parameters
decorator = ast.copy_location(
ast.Call(
func=self._create_module_attribute(
func=_create_module_attribute(
self.DECORATOR_FOR_STATEMENT,
lineno=node.lineno,
col_offset=node.col_offset,
@@ -902,14 +1023,14 @@ class DSLPreprocessor(ast.NodeTransformer):
# else
# return and_(lhs, rhs)
short_circuit_value = ast.Constant(value=False)
helper_func = self._create_module_attribute(
helper_func = _create_module_attribute(
"and_",
top_module_name="cutlass",
use_base_dsl=False,
submodule_name=None,
lineno=node.lineno,
col_offset=node.col_offset,
)
self.import_top_module = True
self.session_data.import_top_module = True
# Transform "or" to "or_"
elif isinstance(node.op, ast.Or):
# Create an if-else statement in AST form
@@ -918,19 +1039,19 @@ class DSLPreprocessor(ast.NodeTransformer):
# else
# return or_(lhs, rhs)
short_circuit_value = ast.Constant(value=True)
helper_func = self._create_module_attribute(
helper_func = _create_module_attribute(
"or_",
top_module_name="cutlass",
use_base_dsl=False,
submodule_name=None,
lineno=node.lineno,
col_offset=node.col_offset,
)
self.import_top_module = True
self.session_data.import_top_module = True
else:
# BoolOp should be either And or Or
raise DSLAstPreprocessorError(
f"Unsupported boolean operation: {node.op}",
filename=self.file_name,
filename=self.session_data.file_name,
snippet=ast.unparse(node),
)
@@ -977,14 +1098,14 @@ class DSLPreprocessor(ast.NodeTransformer):
# Transform "not" to "~" as we overload __invert__
if isinstance(node.op, ast.Not):
func_name = self._create_module_attribute(
func_name = _create_module_attribute(
"not_",
top_module_name="cutlass",
use_base_dsl=False,
submodule_name=None,
lineno=node.lineno,
col_offset=node.col_offset,
)
self.import_top_module = True
self.session_data.import_top_module = True
return ast.copy_location(
ast.Call(func=func_name, args=[node.operand], keywords=[]), node
)
@@ -998,7 +1119,7 @@ class DSLPreprocessor(ast.NodeTransformer):
range_inputs = node.iter.args
check_call = ast.copy_location(
ast.Call(
func=self._create_module_attribute(
func=_create_module_attribute(
"range_value_check", lineno=node.lineno, col_offset=node.col_offset
),
args=range_inputs,
@@ -1021,7 +1142,7 @@ class DSLPreprocessor(ast.NodeTransformer):
"""
check_call = ast.copy_location(
ast.Call(
func=self._create_module_attribute(
func=_create_module_attribute(
"cf_symbol_check", lineno=func.lineno, col_offset=func.col_offset
),
args=[deepcopy(func)],
@@ -1044,18 +1165,18 @@ class DSLPreprocessor(ast.NodeTransformer):
return [check_call, node]
return node
active_symbols = self.scope_manager.get_active_symbols()
active_symbols = self.session_data.scope_manager.get_active_symbols()
with self.scope_manager:
with self.session_data.scope_manager:
if isinstance(node.target, ast.Name):
self.scope_manager.add_to_scope(node.target.id)
self.session_data.scope_manager.add_to_scope(node.target.id)
if range_kind == "range_dynamic":
# Generate a warning
self.issue_deprecation_warning(
message="range_dynamic is deprecated and will be removed in the future, please remove it.",
category=DeprecationWarning,
filename=self.file_name,
filename=self.session_data.file_name,
lineno=node.iter.lineno,
)
@@ -1098,18 +1219,18 @@ class DSLPreprocessor(ast.NodeTransformer):
def _handle_negative_step(self, node, start_expr, stop_expr, step_expr):
# hoist start, stop, step to assignments
start_ori_name = f"start_ori_{self.counter}"
start_ori_name = f"start_ori_{self.session_data.counter}"
start = self._hoist_expr_to_assignments(start_expr, start_ori_name)
stop_ori_name = f"stop_ori_{self.counter}"
stop_ori_name = f"stop_ori_{self.session_data.counter}"
stop = self._hoist_expr_to_assignments(stop_expr, stop_ori_name)
step_ori_name = f"step_ori_{self.counter}"
step_ori_name = f"step_ori_{self.session_data.counter}"
step = self._hoist_expr_to_assignments(step_expr, step_ori_name)
extra_exprs = [start, stop, step]
# Handle possible negative step, generates the following code in Python:
# isNegative = step < 0
isNegative_name = f"isNegative_{self.counter}"
isNegative_name = f"isNegative_{self.session_data.counter}"
isNegative = ast.copy_location(
ast.Assign(
targets=[ast.Name(id=isNegative_name, ctx=ast.Store())],
@@ -1123,7 +1244,7 @@ class DSLPreprocessor(ast.NodeTransformer):
)
# start = stop if isNegative else start
start_name = f"start_{self.counter}"
start_name = f"start_{self.session_data.counter}"
start = self._build_select_and_assign(
name=start_name,
test=ast.Name(id=isNegative_name, ctx=ast.Load()),
@@ -1133,7 +1254,7 @@ class DSLPreprocessor(ast.NodeTransformer):
)
# stop = start if isNegative else stop
stop_name = f"stop_{self.counter}"
stop_name = f"stop_{self.session_data.counter}"
stop = self._build_select_and_assign(
name=stop_name,
test=ast.Name(id=isNegative_name, ctx=ast.Load()),
@@ -1143,7 +1264,7 @@ class DSLPreprocessor(ast.NodeTransformer):
)
# step = -step if isNegative else step
step_name = f"step_{self.counter}"
step_name = f"step_{self.session_data.counter}"
step = self._build_select_and_assign(
name=step_name,
test=ast.Name(id=isNegative_name, ctx=ast.Load()),
@@ -1155,7 +1276,7 @@ class DSLPreprocessor(ast.NodeTransformer):
)
# offset = start + stop if isNegative else 0
offset_name = f"offset_{self.counter}"
offset_name = f"offset_{self.session_data.counter}"
offset = self._build_select_and_assign(
name=offset_name,
test=ast.Name(id=isNegative_name, ctx=ast.Load()),
@@ -1204,7 +1325,7 @@ class DSLPreprocessor(ast.NodeTransformer):
def _create_closure_check_call(self, called_closures, node):
return ast.Expr(
ast.Call(
func=self._create_module_attribute(
func=_create_module_attribute(
"closure_check",
lineno=node.lineno,
col_offset=node.col_offset,
@@ -1225,7 +1346,7 @@ class DSLPreprocessor(ast.NodeTransformer):
if node.orelse:
raise DSLAstPreprocessorError(
"dynamic for loop with else is not supported",
filename=self.file_name,
filename=self.session_data.file_name,
snippet=ast.unparse(node),
)
@@ -1243,7 +1364,7 @@ class DSLPreprocessor(ast.NodeTransformer):
# Add necessary exprs to handle this
if target_var_is_active_before_loop:
# Initialize an extra loop carried variable
loop_carried_var_name = f"loop_carried_var_{self.counter}"
loop_carried_var_name = f"loop_carried_var_{self.session_data.counter}"
pre_loop_expr = ast.copy_location(
ast.Assign(
targets=[ast.Name(id=loop_carried_var_name, ctx=ast.Store())],
@@ -1283,8 +1404,8 @@ class DSLPreprocessor(ast.NodeTransformer):
if called_closures:
exprs.append(self._create_closure_check_call(called_closures, node))
func_name = f"loop_body_{self.counter}"
self.counter += 1
func_name = f"loop_body_{self.session_data.counter}"
self.session_data.counter += 1
func_def = self.create_loop_function(
func_name,
@@ -1329,7 +1450,7 @@ class DSLPreprocessor(ast.NodeTransformer):
# Rewrite to assert_executor(test, msg)
new_node = ast.Expr(
ast.Call(
func=self._create_module_attribute(
func=_create_module_attribute(
self.ASSERT_EXECUTOR, lineno=node.lineno, col_offset=node.col_offset
),
args=[],
@@ -1353,7 +1474,7 @@ class DSLPreprocessor(ast.NodeTransformer):
if func.id == "bool":
return ast.copy_location(
ast.Call(
func=self._create_module_attribute(
func=_create_module_attribute(
self.BOOL_CAST,
lineno=node.lineno,
col_offset=node.col_offset,
@@ -1369,7 +1490,7 @@ class DSLPreprocessor(ast.NodeTransformer):
)
return ast.copy_location(
ast.Call(
func=self._create_module_attribute(
func=_create_module_attribute(
helper_func, lineno=node.lineno, col_offset=node.col_offset
),
args=[node.args[0]],
@@ -1378,12 +1499,12 @@ class DSLPreprocessor(ast.NodeTransformer):
node,
)
elif func.id in ["min", "max"]:
self.import_top_module = True
self.session_data.import_top_module = True
return ast.copy_location(
ast.Call(
func=self._create_module_attribute(
func=_create_module_attribute(
func.id,
top_module_name="cutlass",
use_base_dsl=False,
submodule_name=None,
lineno=node.lineno,
col_offset=node.col_offset,
@@ -1417,7 +1538,7 @@ class DSLPreprocessor(ast.NodeTransformer):
def create_downcast_call(arg):
return ast.copy_location(
ast.Call(
func=self._create_module_attribute(
func=_create_module_attribute(
self.IMPLICIT_DOWNCAST_NUMERIC_TYPE,
submodule_name="typing",
lineno=node.lineno,
@@ -1429,7 +1550,7 @@ class DSLPreprocessor(ast.NodeTransformer):
arg,
)
module = self.function_globals.get(func.value.id)
module = self.session_data.function_globals.get(func.value.id)
if isinstance(module, ModuleType) and module.__package__.endswith(
"._mlir.dialects"
):
@@ -1457,18 +1578,18 @@ class DSLPreprocessor(ast.NodeTransformer):
return node
def visit_ClassDef(self, node):
self.class_name = node.name
self.session_data.class_name = node.name
self.generic_visit(node)
self.class_name = None
self.session_data.class_name = None
return node
def _visit_target(self, target):
if isinstance(target, ast.Name):
self.scope_manager.add_to_scope(target.id)
self.session_data.scope_manager.add_to_scope(target.id)
elif isinstance(target, ast.Tuple):
for t in target.elts:
if isinstance(t, ast.Name):
self.scope_manager.add_to_scope(t.id)
self.session_data.scope_manager.add_to_scope(t.id)
def visit_Assign(self, node):
for target in node.targets:
@@ -1486,7 +1607,7 @@ class DSLPreprocessor(ast.NodeTransformer):
if node.id in ["max", "min", "any", "all"] and isLoad:
return ast.copy_location(
ast.Call(
func=self._create_module_attribute(
func=_create_module_attribute(
"redirect_builtin_function",
lineno=node.lineno,
col_offset=node.col_offset,
@@ -1558,41 +1679,41 @@ class DSLPreprocessor(ast.NodeTransformer):
return new_decorator_list
def visit_FunctionDef(self, node):
with self.scope_manager:
self.function_counter += 1
self.function_name = node.name
if self.function_depth > 0:
self.local_closures.add(node.name)
with self.session_data.scope_manager:
self.session_data.function_counter += 1
self.session_data.function_name = node.name
if self.session_data.function_depth > 0:
self.session_data.local_closures.add(node.name)
self.function_depth += 1
self.session_data.function_depth += 1
# Add function name and arguments
self.scope_manager.add_to_scope(node.name)
self.session_data.scope_manager.add_to_scope(node.name)
for arg in node.args.args:
self.scope_manager.add_to_scope(arg.arg)
self.session_data.scope_manager.add_to_scope(arg.arg)
arg.annotation = None
for arg in node.args.kwonlyargs:
self.scope_manager.add_to_scope(arg.arg)
self.session_data.scope_manager.add_to_scope(arg.arg)
arg.annotation = None
for arg in node.args.posonlyargs:
self.scope_manager.add_to_scope(arg.arg)
self.session_data.scope_manager.add_to_scope(arg.arg)
arg.annotation = None
self.generic_visit(node)
self.function_depth -= 1
self.session_data.function_depth -= 1
# Remove .jit and .kernel decorators
node.decorator_list = self.remove_dsl_decorator(node.decorator_list)
return node
def visit_With(self, node):
with self.scope_manager:
with self.session_data.scope_manager:
for item in node.items:
if isinstance(item.optional_vars, ast.Name):
self.scope_manager.add_to_scope(item.optional_vars.id)
self.session_data.scope_manager.add_to_scope(item.optional_vars.id)
self.generic_visit(node)
return node
@@ -1604,8 +1725,8 @@ class DSLPreprocessor(ast.NodeTransformer):
check = self._insert_cf_symbol_check(node.test.func)
return [check, node]
active_symbols = self.scope_manager.get_active_symbols()
with self.scope_manager:
active_symbols = self.session_data.scope_manager.get_active_symbols()
with self.session_data.scope_manager:
# Check for early exit and raise exception
self.check_early_exit(node, "while")
@@ -1616,8 +1737,8 @@ class DSLPreprocessor(ast.NodeTransformer):
if called_closures:
exprs.append(self._create_closure_check_call(called_closures, node))
func_name = f"while_region_{self.counter}"
self.counter += 1
func_name = f"while_region_{self.session_data.counter}"
self.session_data.counter += 1
func_def = self.create_while_function(
func_name, node, write_args, full_write_args_count
@@ -1627,14 +1748,14 @@ class DSLPreprocessor(ast.NodeTransformer):
return exprs + [func_def] + assign
def visit_Try(self, node):
with self.scope_manager:
with self.session_data.scope_manager:
self.generic_visit(node)
return node
def visit_ExceptHandler(self, node):
with self.scope_manager:
with self.session_data.scope_manager:
if node.name: # Exception variable
self.scope_manager.add_to_scope(node.name)
self.session_data.scope_manager.add_to_scope(node.name)
self.generic_visit(node)
return node
@@ -1671,7 +1792,7 @@ class DSLPreprocessor(ast.NodeTransformer):
if has_self:
fix_self = ast.Expr(
value=ast.Call(
func=self._create_module_attribute(
func=_create_module_attribute(
"copy_members", lineno=node.lineno, col_offset=node.col_offset
),
args=[
@@ -1694,7 +1815,7 @@ class DSLPreprocessor(ast.NodeTransformer):
# Emit
# node if type(pred) == bool else select_(pred, body, orelse)
# so if pred is a python bool, use python to short-circuit and avoid emit arith.select
self.import_top_module = True
self.session_data.import_top_module = True
return ast.copy_location(
ast.IfExp(
test=ast.Compare(
@@ -1708,8 +1829,8 @@ class DSLPreprocessor(ast.NodeTransformer):
),
body=node, # Original ternary expression
orelse=ast.Call(
func=self._create_module_attribute(
"select_", top_module_name="cutlass", submodule_name=None
func=_create_module_attribute(
"select_", use_base_dsl=False, submodule_name=None
),
args=[
node.test,
@@ -1756,7 +1877,7 @@ class DSLPreprocessor(ast.NodeTransformer):
call = ast.copy_location(
ast.Call(
func=self._create_module_attribute(self.COMPARE_EXECUTOR),
func=_create_module_attribute(self.COMPARE_EXECUTOR),
args=[],
keywords=keywords,
),
@@ -1772,8 +1893,8 @@ class DSLPreprocessor(ast.NodeTransformer):
check = self._insert_cf_symbol_check(node.test.func)
return [check, node]
active_symbols = self.scope_manager.get_active_symbols()
with self.scope_manager:
active_symbols = self.session_data.scope_manager.get_active_symbols()
with self.session_data.scope_manager:
# Check for early exit and raise exception
self.check_early_exit(node, "if")
@@ -1784,8 +1905,8 @@ class DSLPreprocessor(ast.NodeTransformer):
if called_closures:
exprs.append(self._create_closure_check_call(called_closures, node))
func_name = f"if_region_{self.counter}"
self.counter += 1
func_name = f"if_region_{self.session_data.counter}"
self.session_data.counter += 1
func_def = self.create_if_function(
func_name, node, yield_args, full_write_args_count
@@ -1796,7 +1917,7 @@ class DSLPreprocessor(ast.NodeTransformer):
def generate_get_locals_or_none_call(self, write_args):
return ast.Call(
func=self._create_module_attribute("get_locals_or_none"),
func=_create_module_attribute("get_locals_or_none"),
args=[
ast.Call(
func=ast.Name(id="locals", ctx=ast.Load()), args=[], keywords=[]
@@ -1842,10 +1963,10 @@ class DSLPreprocessor(ast.NodeTransformer):
defaults=[],
)
then_block_name = f"then_block_{self.counter}"
else_block_name = f"else_block_{self.counter}"
elif_region_name = f"elif_region_{self.counter}"
self.counter += 1
then_block_name = f"then_block_{self.session_data.counter}"
else_block_name = f"else_block_{self.session_data.counter}"
elif_region_name = f"elif_region_{self.session_data.counter}"
self.session_data.counter += 1
# Create then block
then_block = ast.copy_location(
@@ -1872,7 +1993,7 @@ class DSLPreprocessor(ast.NodeTransformer):
# Create decorator
decorator = ast.copy_location(
ast.Call(
func=self._create_module_attribute(
func=_create_module_attribute(
self.DECORATOR_IF_STATEMENT,
lineno=node.lineno,
col_offset=node.col_offset,
@@ -1914,7 +2035,7 @@ class DSLPreprocessor(ast.NodeTransformer):
# No write_args case - only then_block needed
execute_call = ast.copy_location(
ast.Call(
func=self._create_module_attribute(
func=_create_module_attribute(
self.IF_EXECUTOR, lineno=node.lineno, col_offset=node.col_offset
),
args=[],
@@ -2001,7 +2122,7 @@ class DSLPreprocessor(ast.NodeTransformer):
execute_call = ast.copy_location(
ast.Call(
func=self._create_module_attribute(
func=_create_module_attribute(
self.IF_EXECUTOR, lineno=node.lineno, col_offset=node.col_offset
),
args=[],
@@ -2072,7 +2193,7 @@ class DSLPreprocessor(ast.NodeTransformer):
]
decorator = ast.copy_location(
ast.Call(
func=self._create_module_attribute(
func=_create_module_attribute(
self.DECORATOR_WHILE_STATEMENT,
lineno=node.lineno,
col_offset=node.col_offset,
@@ -2084,9 +2205,9 @@ class DSLPreprocessor(ast.NodeTransformer):
)
# Section: Shared initialization for before and after blocks
while_before_block_name = f"while_before_block_{self.counter}"
while_after_block_name = f"while_after_block_{self.counter}"
self.counter += 1
while_before_block_name = f"while_before_block_{self.session_data.counter}"
while_after_block_name = f"while_after_block_{self.session_data.counter}"
self.session_data.counter += 1
block_args_args = [ast.arg(arg=var, annotation=None) for var in write_args]
block_args = ast.arguments(
posonlyargs=[],
@@ -2169,7 +2290,7 @@ class DSLPreprocessor(ast.NodeTransformer):
]
execute_call = ast.Call(
func=self._create_module_attribute(
func=_create_module_attribute(
self.WHILE_EXECUTOR, lineno=node.lineno, col_offset=node.col_offset
),
args=[],
python/CuTeDSL/cutlass/base_dsl/cache_helpers.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/common.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/compiler.py
+10 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
@@ -626,6 +626,10 @@ class CompileCallable:
while hasattr(func, "__wrapped__"):
func = func.__wrapped__
from .dsl import BaseDSL
BaseDSL._lazy_initialize_dsl(func)
if not hasattr(func, "_dsl_object"):
raise DSLRuntimeError(
f"Function {func} is not decorated with jit decorator."
@@ -641,4 +645,9 @@ class CompileCallable:
else:
compile_options = self._compile_options
func._dsl_object.compile_options = compile_options
# Preprocess the function if not already preprocessed
func._dsl_object._preprocess_and_replace_code(func)
# Run the function
return func._dsl_object._func(func, *args, **kwargs)
python/CuTeDSL/cutlass/base_dsl/dsl.py
+86 -127
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
@@ -229,38 +229,6 @@ def new_from_mlir_values(obj, values):
assert len(values) == 0, f"{obj} expects 0 values, but got {values}"
return obj
@dataclass(frozen=True)
class DSLLocation:
"""
Represents Python source location information for MLIR DSL code.
Attributes:
filename (str): Name of the Python source file.
lineno (int): Line number in the source file.
col_offset (int): Column offset in the source line.
function_name (str): Name of the function in which the location occurs.
This is used primarily to annotate or trace locations in generated MLIR IR
back to the original Python code for better diagnostic and debugging.
"""
filename: str
lineno: int
col_offset: int
function_name: str
@dataclass
class PreprocessSessionData:
"""
Holds metadata and transformed AST related to a DSL preprocessing session.
Attributes:
decorator_globals (dict): The global variables from the decorator's environment,
captured for possible AST or code evaluation during preprocessing.
"""
decorator_globals: dict
class DSLSingletonMeta(type):
"""
@@ -314,6 +282,27 @@ class DSLSingletonMeta(type):
log().info(f"DSL singleton instances after clearing: {cls._instances}")
@dataclass(frozen=True)
class DSLLocation:
"""
Represents Python source location information for MLIR DSL code.
Attributes:
filename (str): Name of the Python source file.
lineno (int): Line number in the source file.
col_offset (int): Column offset in the source line.
function_name (str): Name of the function in which the location occurs.
This is used primarily to annotate or trace locations in generated MLIR IR
back to the original Python code for better diagnostic and debugging.
"""
filename: str
lineno: int
col_offset: int
function_name: str
class BaseDSL(metaclass=DSLSingletonMeta):
gpu_module = None
_env_class = EnvironmentVarManager
@@ -354,7 +343,6 @@ class BaseDSL(metaclass=DSLSingletonMeta):
self.name = name
self.compiler_provider = compiler_provider
self.pass_sm_arch_name = pass_sm_arch_name
self.preprocess_session_data = None
self.decorator_location = None
self.no_cache = False
self.device_compilation_only = device_compilation_only
@@ -437,44 +425,40 @@ class BaseDSL(metaclass=DSLSingletonMeta):
return dkwargs.pop("preprocess", True)
@staticmethod
def _get_original_function(fcn_ptr, name):
def _lazy_initialize_dsl(func):
"""
Get the original function from the decorated function
Lazy initialization of DSL object if has not been initialized
"""
while not hasattr(fcn_ptr, "__name__") or fcn_ptr.__name__ != name:
# If the function is wrapped with functools, get from __wrapped__
if hasattr(fcn_ptr, "__wrapped__"):
fcn_ptr = fcn_ptr.__wrapped__
elif isinstance(fcn_ptr, staticmethod):
fcn_ptr = fcn_ptr.__func__
# If the function is wrapped manually, it's the first in clousure
elif callable(fcn_ptr.__closure__[0].cell_contents):
fcn_ptr = fcn_ptr.__closure__[0].cell_contents
else:
raise DSLRuntimeError(
f"Cannot find the original function {name} in the closure chain"
)
return fcn_ptr
if hasattr(func, "_dsl_cls"):
func._dsl_object = func._dsl_cls._get_dsl()
delattr(func, "_dsl_cls")
@staticmethod
def _preprocess_and_replace_code(func):
"""
Run ast transformation and return the materialized function pointer
"""
# Ensure the DSL instance is materialized before touching _dsl_object
BaseDSL._lazy_initialize_dsl(func)
if hasattr(func, "_preprocess_session_data"):
# If the function ptr is already materialized, use the existing one
func._dsl_object.preprocess_session_data = func._preprocess_session_data
func._dsl_object.decorator_location = func._decorator_location
transformed_ast = func._dsl_object.run_preprocessor(func)
fcn_ptr = func._dsl_object.get_function_ptr(func, transformed_ast)
# Update the decorator location to the new function
func._dsl_object.decorator_location = func._decorator_location
if getattr(func, "_preprocessed", False) is True:
# already preprocessed, skip
return
if not func._dsl_object.enable_preprocessor:
func._preprocessed = True
return
fcn_ptr = func._dsl_object.run_preprocessor(func)
if fcn_ptr:
func.__code__ = (
fcn_ptr.__code__
if not isinstance(fcn_ptr, staticmethod)
else fcn_ptr.__func__.__code__
)
return func
@staticmethod
def jit_runner(cls, executor_name, frame, *dargs, **dkwargs):
@@ -485,22 +469,26 @@ class BaseDSL(metaclass=DSLSingletonMeta):
def jit_runner_decorator(func):
# Run preprocessor that alters AST
func._dsl_object = cls._get_dsl()
func._dsl_cls = cls
func._decorator_location = BaseDSL.get_location_from_frame(frame)
if (
func._dsl_object.enable_preprocessor
and func._dsl_object._can_preprocess(**dkwargs)
if not hasattr(func, "_preprocessed") and not BaseDSL._can_preprocess(
**dkwargs
):
# For an annotated function, add some DSL attributes
# When materializing the AST, we need decorator's frame
func._preprocess_session_data = PreprocessSessionData(
decorator_globals=frame.f_globals,
)
BaseDSL._preprocess_and_replace_code(func)
func._preprocessed = True
@wraps(func)
def jit_wrapper(*args, **kwargs):
return getattr(func._dsl_object, executor_name)(func, *args, **kwargs)
BaseDSL._preprocess_and_replace_code(func)
custom_name = getattr(jit_wrapper, "_name_prefix", None)
if custom_name:
return getattr(func._dsl_object, executor_name)(
func, *args, **kwargs, _name_prefix=custom_name
)
else:
return getattr(func._dsl_object, executor_name)(
func, *args, **kwargs
)
def set_name_prefix(name: str):
jit_wrapper._name_prefix = name
@@ -702,26 +690,6 @@ class BaseDSL(metaclass=DSLSingletonMeta):
"""
pass
def _get_globals(self):
"""
Combines global and local variables from the current context and the
caller's frame comes. This includes the current module's globals, the
global variables from the caller's frame, and the local variables from
the caller's frame.
"self.frame" is used to fetch the caller's frame.
AST preprocessor generates a new python code, so the resulting globals
dictionary is used to execute the python code.
"""
all_globals = {}
if (
self.preprocess_session_data
and self.preprocess_session_data.decorator_globals
):
all_globals.update(self.preprocess_session_data.decorator_globals)
return all_globals
@abstractmethod
def _is_tensor_descriptor(self, maybe_tensor_descriptor) -> bool:
pass
@@ -986,18 +954,11 @@ class BaseDSL(metaclass=DSLSingletonMeta):
@staticmethod
def get_location_from_frame(frame):
frameInfo = inspect.getframeinfo(frame)
return DSLLocation(
filename=frameInfo.filename,
lineno=frameInfo.lineno,
col_offset=(
frameInfo.positions.col_offset if hasattr(frameInfo, "positions") else 0
),
function_name=(
"".join([c.strip() for c in frameInfo.code_context])
if frameInfo.code_context
else frameInfo.function
),
filename=inspect.getsourcefile(frame),
lineno=frame.f_lineno,
col_offset=0,
function_name=frame.f_code.co_name,
)
def get_ir_location(self, location: DSLLocation = None):
@@ -1352,8 +1313,6 @@ class BaseDSL(metaclass=DSLSingletonMeta):
self.num_kernels = 0
# reset the compile options after the compilation is done.
self.compile_options = CompileOptions()
# reset preprocess session data after the compilation is done.
self.preprocess_session_data = None
# reset decorator location after the compilation is done.
self.decorator_location = None
@@ -1468,36 +1427,37 @@ class BaseDSL(metaclass=DSLSingletonMeta):
return result
def run_preprocessor(self, funcBody):
if not hasattr(funcBody, "_preprocessed"):
function_name = funcBody.__name__
self.funcBody = funcBody
log().info("Started preprocessing [%s]", function_name)
exec_globals = self._get_globals()
transformed_ast = self.preprocessor.transform(funcBody, exec_globals)
def run_preprocessor(self, original_function):
function_name = original_function.__name__
self.funcBody = original_function
log().info("Started preprocessing [%s]", function_name)
exec_globals = {}
if original_function.__globals__ is not None:
exec_globals.update(original_function.__globals__)
with self.preprocessor.get_session() as preprocessor_session:
transformed_ast = preprocessor_session.transform(
original_function, exec_globals
)
if self.envar.print_after_preprocessor:
log().info(
f"# Printing unparsed AST after preprocess of func=`{function_name}` id=`{id(funcBody)}`"
f"# Printing unparsed AST after preprocess of func=`{function_name}` id=`{id(original_function)}`"
)
DSLPreprocessor.print_ast(transformed_ast)
funcBody._preprocessed = True
return transformed_ast
return None
file_name = inspect.getsourcefile(original_function)
code_object = compile(
transformed_ast,
filename=file_name,
mode="exec",
)
def get_function_ptr(self, original_function, transformed_ast):
file_name = inspect.getsourcefile(original_function)
code_object = compile(
transformed_ast,
filename=file_name,
mode="exec",
)
original_function._preprocessed = True
return self.preprocessor.exec(
original_function.__name__,
original_function,
code_object,
self._get_globals(),
)
return preprocessor_session.exec(
original_function.__name__,
original_function,
code_object,
exec_globals,
)
def _get_function_bound_args(self, sig, func_name, *args, **kwargs):
"""
@@ -1573,7 +1533,6 @@ class BaseDSL(metaclass=DSLSingletonMeta):
pipeline = kwargs.pop("pipeline", None)
gpu_module_attrs = kwargs.pop("gpu_module_attrs", {})
self.decorator_location = getattr(funcBody, "_decorator_location", None)
# Disable cache
no_cache = kwargs.pop("no_cache", False)
python/CuTeDSL/cutlass/base_dsl/env_manager.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/export/__init__.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/export/c_header_generator.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/export/export.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/jit_executor.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/runtime/__init__.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/runtime/cuda.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/runtime/device_tensor.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/runtime/dlpack_types.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/runtime/jit_arg_adapters.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/runtime/stream_adapter.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/runtime/tensor_descriptor.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/tvm_ffi_builder/__init__.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/tvm_ffi_builder/call_provider.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/tvm_ffi_builder/mlir_builder.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/tvm_ffi_builder/spec.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/tvm_ffi_builder/tvm_ffi_builder.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/typing.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/utils/__init__.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/utils/logger.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/utils/stacktrace.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/base_dsl/utils/timer.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/__init__.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/_tvm_ffi_args_spec_converter.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/algorithm.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/arch/__init__.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/arch/elect.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/arch/mbar.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/arch/numeric_conversion.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/arch/nvvm_wrappers.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/arch/smem.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/arch/tmem.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/atom.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/core.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/export/__init__.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/export/c_header_generator.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/math.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/nvgpu/__init__.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/nvgpu/common.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/nvgpu/cpasync/__init__.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/nvgpu/cpasync/copy.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/nvgpu/cpasync/helpers.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/nvgpu/helpers.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/nvgpu/tcgen05/__init__.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/nvgpu/tcgen05/copy.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/nvgpu/tcgen05/helpers.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/nvgpu/tcgen05/mma.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/nvgpu/warp/__init__.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/nvgpu/warp/copy.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/nvgpu/warp/mma.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/nvgpu/warpgroup/__init__.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/nvgpu/warpgroup/helpers.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/nvgpu/warpgroup/mma.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/runtime.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/tensor.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/testing.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/tuple.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cute/typing.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cutlass_dsl/__init__.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cutlass_dsl/cuda_jit_executor.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cutlass_dsl/cuda_stream_adapter.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cutlass_dsl/cutlass.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cutlass_dsl/cutlass_ast_decorators.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cutlass_dsl/tree_utils.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/cutlass_dsl/tvm_ffi_provider.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/impl_utils.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/pipeline/__init__.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/pipeline/helpers.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/pipeline/sm100.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/pipeline/sm90.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/torch.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/utils/__init__.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/utils/blackwell_helpers.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/utils/blockscaled_layout.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/utils/distributed.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/utils/grouped_gemm_tile_scheduler_helper.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/utils/hardware_info.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/utils/hopper_helpers.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/utils/layout.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/utils/mixed_input_helpers.py
+1 -1
View File
@@ -1,4 +1,4 @@
# Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: BSD-3-Clause
# Redistribution and use in source and binary forms, with or without
python/CuTeDSL/cutlass/utils/print_latex.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/utils/smem_allocator.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/utils/static_persistent_tile_scheduler.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/utils/tensormap_manager.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/cutlass/utils/tmem_allocator.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# Use of this software is governed by the terms and conditions of the
python/CuTeDSL/prep_editable_install.py
+1 -1
View File
@@ -1,4 +1,4 @@
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-FileCopyrightText: Copyright (c) 2025 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: LicenseRef-NvidiaProprietary
#
# NVIDIA CORPORATION, its affiliates and licensors retain all intellectual

Some files were not shown because too many files have changed in this diff Show More