Skip to content

Commit 351a431

Browse files
committed
support nvshmem
1 parent b4b2684 commit 351a431

22 files changed

Lines changed: 2026 additions & 3 deletions

File tree

include/triton/Conversion/TritonGPUToLLVM/PatternTritonGPUOpToLLVM.h

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -102,6 +102,11 @@ void populatePrintOpToLLVMPattern(LLVMTypeConverter &typeConverter,
102102
const TargetInfoBase &targetInfo,
103103
PatternBenefit benefit);
104104

105+
void populateExternCallOpToLLVMPattern(LLVMTypeConverter &typeConverter,
106+
RewritePatternSet &patterns,
107+
const TargetInfoBase &targetInfo,
108+
PatternBenefit benefit);
109+
105110
void populateInstrumentationToLLVMPatterns(LLVMTypeConverter &typeConverter,
106111
const TargetInfoBase &targetInfo,
107112
RewritePatternSet &patterns,

include/triton/Dialect/Triton/IR/TritonOps.td

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -845,6 +845,32 @@ def TT_MapElementwiseReturnOp: TT_Op<"map_elementwise.return",
845845
let assemblyFormat = "attr-dict ($result^ `:` type($result))?";
846846
}
847847

848+
//
849+
// External Call op
850+
//
851+
def TT_ExternCallOp : TT_Op<"extern_call", [
852+
DeclareOpInterfaceMethods<MemoryEffectsOpInterface>,
853+
ConditionallySpeculatable,
854+
]> {
855+
856+
let description = [{
857+
call an external function $symbol implemented in $libpath/$libname with $args
858+
return $libpath/$libname:$symbol($args...)
859+
}];
860+
861+
let arguments = (ins Variadic<TT_Type>:$srcs, StrAttr:$libname, StrAttr:$libpath, StrAttr:$symbol, BoolAttr:$pure);
862+
863+
let results = (outs Variadic<TT_Type>:$result);
864+
865+
let assemblyFormat = "operands attr-dict `:` functional-type(operands, $result)";
866+
867+
let extraClassDeclaration = [{
868+
// Interface method for ConditionallySpeculatable.
869+
Speculation::Speculatability getSpeculatability();
870+
}];
871+
872+
}
873+
848874
//
849875
// External Elementwise op
850876
//

lib/Conversion/TritonGPUToLLVM/CMakeLists.txt

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -21,6 +21,7 @@ add_triton_library(TritonGPUToLLVM
2121
TypeConverter.cpp
2222
Utility.cpp
2323
ViewOpToLLVM.cpp
24+
ExternCallOpToLLVM.cpp
2425

2526
DEPENDS
2627
TritonGPUConversionPassIncGen
Lines changed: 61 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,61 @@
1+
#include "mlir/Conversion/LLVMCommon/Pattern.h"
2+
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
3+
#include "mlir/Support/LLVM.h"
4+
#include "triton/Conversion/TritonGPUToLLVM/ElementwiseOpToLLVMBase.h"
5+
#include "triton/Conversion/TritonGPUToLLVM/PatternTritonGPUOpToLLVM.h"
6+
#include "triton/Conversion/TritonGPUToLLVM/TargetInfoBase.h"
7+
#include "triton/Conversion/TritonGPUToLLVM/Utility.h"
8+
#include "triton/Dialect/TritonGPU/IR/Dialect.h"
9+
10+
namespace {
11+
12+
class ExternCallOpConversion
13+
: public ConvertOpToLLVMPattern<triton::ExternCallOp> {
14+
public:
15+
ExternCallOpConversion(const LLVMTypeConverter &converter,
16+
const PatternBenefit &benefit)
17+
: ConvertOpToLLVMPattern<triton::ExternCallOp>(converter, benefit) {}
18+
19+
LogicalResult
20+
matchAndRewrite(triton::ExternCallOp op, OpAdaptor adaptor,
21+
ConversionPatternRewriter &rewriter) const override {
22+
auto loc = op->getLoc();
23+
24+
if (op->getNumResults() > 1) {
25+
llvm::errs() << "ExternCallConversion does not support multi outs.";
26+
return failure();
27+
}
28+
29+
LLVM::LLVMVoidType voidTy = void_ty(op->getContext());
30+
auto newOperands = adaptor.getOperands();
31+
Type retType =
32+
op->getNumResults() == 0
33+
? voidTy
34+
: this->getTypeConverter()->convertType(op->getResult(0).getType());
35+
std::string funcName = op.getSymbol().str();
36+
StringRef libname = op.getLibname();
37+
StringRef libpath = op.getLibpath();
38+
39+
Operation *externCallOp;
40+
Type funcType = mlir::triton::gpu::getFunctionType(retType, newOperands);
41+
LLVM::LLVMFuncOp funcOp = mlir::triton::gpu::appendOrGetExternFuncOp(
42+
rewriter, op, funcName, funcType, libname, libpath);
43+
externCallOp = LLVM::createLLVMCallOp(rewriter, loc, funcOp, newOperands);
44+
45+
if (op->getNumResults() == 0) {
46+
rewriter.eraseOp(op);
47+
} else {
48+
rewriter.replaceOp(op, externCallOp->getResult(0));
49+
}
50+
51+
return success();
52+
}
53+
};
54+
55+
} // namespace
56+
57+
void mlir::triton::populateExternCallOpToLLVMPattern(
58+
LLVMTypeConverter &typeConverter, RewritePatternSet &patterns,
59+
const TargetInfoBase &targetInfo, PatternBenefit benefit) {
60+
patterns.add<ExternCallOpConversion>(typeConverter, benefit);
61+
}

lib/Conversion/TritonToTritonGPU/TritonToTritonGPUPass.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -618,6 +618,7 @@ void populateTritonPatterns(TritonGPUTypeConverter &typeConverter,
618618
GenericOpPattern<triton::HistogramOp>,
619619
GenericOpPattern<triton::GatherOp>,
620620
GenericOpPattern<triton::ExternElementwiseOp>,
621+
GenericOpPattern<triton::ExternCallOp>,
621622
GenericOpPattern<triton::PrintOp>,
622623
GenericOpPattern<triton::AssertOp>,
623624
GenericOpPattern<triton::AtomicCASOp>,

lib/Dialect/Triton/IR/Ops.cpp

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1313,6 +1313,24 @@ Speculation::Speculatability ExternElementwiseOp::getSpeculatability() {
13131313
return Speculation::NotSpeculatable;
13141314
}
13151315

1316+
// -- ExternCallOp --
1317+
void ExternCallOp::getEffects(
1318+
SmallVectorImpl<SideEffects::EffectInstance<MemoryEffects::Effect>>
1319+
&effects) {
1320+
if (getPure())
1321+
return;
1322+
effects.emplace_back(MemoryEffects::Write::get(),
1323+
SideEffects::DefaultResource::get());
1324+
effects.emplace_back(MemoryEffects::Read::get(),
1325+
SideEffects::DefaultResource::get());
1326+
}
1327+
1328+
Speculation::Speculatability ExternCallOp::getSpeculatability() {
1329+
if (getPure())
1330+
return Speculation::Speculatable;
1331+
return Speculation::NotSpeculatable;
1332+
}
1333+
13161334
// -- GatherOp --
13171335
LogicalResult GatherOp::verify() {
13181336
RankedTensorType indicesTy = getIndices().getType();

python/src/ir.cc

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1673,6 +1673,14 @@ void init_triton_ir(py::module &&m) {
16731673
return self.create<ExternElementwiseOp>(retType, argList, libName,
16741674
libPath, symbol, isPure);
16751675
})
1676+
.def("create_extern_call",
1677+
[](TritonOpBuilder &self, const std::string &libName,
1678+
const std::string &libPath, const std::string &symbol,
1679+
std::vector<Value> &argList, const std::vector<Type> &retTypes,
1680+
bool isPure) -> OpState {
1681+
return self.create<ExternCallOp>(retTypes, argList, libName,
1682+
libPath, symbol, isPure);
1683+
})
16761684
// Built-in instruction
16771685
.def("create_get_program_id",
16781686
[](TritonOpBuilder &self, int axis) -> Value {

python/triton/experimental/tle/language/raw/core.py

Lines changed: 29 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,26 @@
11
import triton.language as tl
22
from triton.language.core import builtin, constexpr as tl_constexpr, tensor
33
from triton.experimental.tle.language.gpu import buffered_tensor
4+
import importlib.util
5+
6+
7+
def _pointer_type_hash(self):
8+
return hash((self.name, self.element_ty, "tt_ptr"))
9+
10+
11+
def patch_hash_method_for_pointer_type():
12+
elem_dtype_list = tl.core.dtype.SINT_TYPES + tl.core.dtype.UINT_TYPES + tl.core.dtype.FP_TYPES + tl.core.dtype.OTHER_TYPES
13+
for elem_dtype in elem_dtype_list:
14+
ptr_ty = type(tl.core.pointer_type(tl.core.dtype(elem_dtype)))
15+
ptr_ty.__hash__ = _pointer_type_hash
16+
17+
18+
def import_from_path(file_path):
19+
module_name = f"_imported_{abs(hash(file_path))}"
20+
spec = importlib.util.spec_from_file_location(module_name, file_path)
21+
module = importlib.util.module_from_spec(spec)
22+
spec.loader.exec_module(module)
23+
return module
424

525

626
def _resolve_alias_indices(func, llvm, handles, output_indices, _semantic):
@@ -42,14 +62,22 @@ def _normalize_hint(hint):
4262

4363
def _tle_raw_call(func, args, *, output_indices, hint, smem, _semantic):
4464
hint = _normalize_hint(hint)
45-
handles = [arg.handle for arg in args]
4665
if getattr(func, "deferred", False):
66+
handles = [arg.handle for arg in args]
4767
if output_indices is None:
4868
raise RuntimeError("deferred tle_raw.call requires explicit output_indices=")
4969
alias_indices = output_indices
5070
source_id = func.register_pending_source(hint=hint)
5171
dsl_region_op = func.create_region_deferred(_semantic.builder, source_id, handles, alias_indices, hint)
5272
else:
73+
if func.compiler.lower() == "nvcc" or (func.compiler.lower() == "clang" and func.target.lower() == "bc"):
74+
patch_hash_method_for_pointer_type()
75+
module = import_from_path(func.extern_file)
76+
target_fn = getattr(module, func.extern_func_name)
77+
ret = target_fn(*args, _semantic=_semantic)
78+
return ret
79+
80+
handles = [arg.handle for arg in args]
5381
context = _semantic.builder.get_context()
5482
llvm = func.make_llvm(context)
5583
alias_indices = _resolve_alias_indices(func, llvm, handles, output_indices, _semantic)

python/triton/experimental/tle/raw/cuda/runtime.py

Lines changed: 139 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -8,6 +8,11 @@
88
from typing import Any, Final
99

1010
import torch
11+
import tempfile
12+
import signal
13+
from triton import knobs
14+
from triton.runtime.errors import PTXASError
15+
from functools import partial
1116

1217
from triton._C.libtriton import llvm # pyright: ignore[reportMissingImports]
1318
from triton._C.libtriton.tle.llvm import parse_llvm_ir # pyright: ignore[reportMissingImports]
@@ -17,6 +22,16 @@
1722
CLANG = os.getenv("CLANG", "clang")
1823
CLANG_FLAGS = shlex.split(os.getenv("CLANG_FLAGS", ""))
1924

25+
NVCC = os.getenv("NVCC", "nvcc")
26+
NVCC_FLAGS = shlex.split(os.getenv("NVCC_FLAGS", ""))
27+
28+
PTXAS = os.getenv("PTXAS", "ptxas")
29+
30+
NVLINK = os.getenv("NVLINK", "nvlink")
31+
NVLINK_FLAGS = shlex.split(os.getenv("NVLINK_FLAGS", ""))
32+
33+
OPT = os.getenv("OPT", "opt")
34+
2035

2136
def _sanitize_clang_ir(ir: str) -> str:
2237
# Newer clang emits attributes that this Triton branch's LLVM parser does
@@ -43,23 +58,113 @@ def _get_cuda_gpu_arch() -> str:
4358
if arch:
4459
return f"--cuda-gpu-arch={arch}"
4560
major, minor = torch.cuda.get_device_capability()
46-
return f"--cuda-gpu-arch=sm_{major}{minor}"
61+
suffix = "a" if major >= 9 else ""
62+
return f"--cuda-gpu-arch=sm_{major}{minor}{suffix}"
63+
64+
65+
def make_cubin_inspection_hook(cuda_self, triton_self, stages, options, language, capability):
66+
67+
def make_cubin(self, src, metadata, opt, capability):
68+
fsrc_cuda = cuda_self.source_file
69+
arch = _get_cuda_gpu_arch().split('=')[1]
70+
fbin_cuda = tempfile.NamedTemporaryFile(delete=False, suffix='.o').name
71+
72+
build = subprocess.run([NVCC, "-c", "-rdc=true", f"-arch={arch}", *NVCC_FLAGS, "-o", fbin_cuda, fsrc_cuda],
73+
capture_output=True)
74+
assert build.returncode == 0, (f"nvcc failed\nstderr:\n{build.stderr.decode()}")
75+
76+
with tempfile.NamedTemporaryFile(delete=False, mode='w', suffix='.ptx') as fsrc_triton, \
77+
tempfile.NamedTemporaryFile(delete=False, mode='r', suffix='.log') as flog:
78+
fsrc_triton.write(src)
79+
fsrc_triton.flush()
80+
fbin_triton = fsrc_triton.name + '.o'
81+
fbin_combined = fbin_triton + '.combined.cubin'
82+
83+
compile_only_cmds = ['-c']
84+
line_info = ["-lineinfo", "-suppress-debug-info"] if knobs.compilation.disable_line_info else ["-lineinfo"]
85+
fmad = [] if opt.enable_fp_fusion else ['--fmad=false']
86+
disable_opt = ['--opt-level', '0'] if knobs.nvidia.disable_ptxas_opt else []
87+
ptx_extra_options = opt.ptx_options.split(" ") if opt.ptx_options else []
88+
89+
ptxas_cmd = [
90+
PTXAS, *compile_only_cmds, *line_info, *fmad, '-v', *disable_opt, *ptx_extra_options,
91+
f'--gpu-name={arch}', fsrc_triton.name, '-o', fbin_triton
92+
]
93+
94+
try:
95+
subprocess.run(ptxas_cmd, check=True, close_fds=False, stderr=flog)
96+
if os.path.exists(fsrc_triton.name):
97+
os.remove(fsrc_triton.name)
98+
if os.path.exists(flog.name):
99+
os.remove(flog.name)
100+
except subprocess.CalledProcessError as e:
101+
with open(flog.name) as log_file:
102+
log = log_file.read()
103+
if os.path.exists(flog.name):
104+
os.remove(flog.name)
105+
106+
if e.returncode == 255:
107+
error = 'Internal Triton PTX codegen error'
108+
elif e.returncode == 128 + signal.SIGSEGV:
109+
error = '`ptxas` raised SIGSEGV'
110+
else:
111+
error = f'`ptxas` failed with error code {e.returncode}'
112+
raise PTXASError(f"{error}\n"
113+
f"`ptxas` stderr:\n{log}\n"
114+
f'Repro command: {" ".join(ptxas_cmd)}\n')
115+
116+
nvlink_cmds = [
117+
NVLINK,
118+
f"-arch={arch}",
119+
*NVLINK_FLAGS,
120+
fbin_triton,
121+
fbin_cuda,
122+
"-o",
123+
fbin_combined,
124+
]
125+
126+
try:
127+
subprocess.run(nvlink_cmds, check=True, close_fds=False, stderr=flog)
128+
except Exception as e:
129+
import logging
130+
logging.error(f"error runing nvlink: {shlex.join(nvlink_cmds)}")
131+
logging.exception(e)
132+
133+
with open(fbin_combined, 'rb') as f:
134+
cubin = f.read()
135+
if os.path.exists(fbin_combined):
136+
os.remove(fbin_combined)
137+
if os.path.exists(fbin_triton):
138+
os.remove(fbin_triton)
139+
if os.path.exists(fbin_cuda):
140+
os.remove(fbin_cuda)
141+
return cubin
142+
143+
stages["cubin"] = lambda src, metadata: make_cubin(triton_self, src, metadata, options, triton_self.target.arch)
47144

48145

49146
class CUDAJITFunction(object):
50147

51148
def __init__(self, fn: Any, file: Path, *args, **kwargs) -> None:
52-
super().__init__(*args, **{k: v for k, v in kwargs.items() if k not in ("extern_func_name", "deferred")})
149+
super().__init__()
150+
# super().__init__(*args, **{k: v for k, v in kwargs.items() if k not in ("extern_func_name", "deferred")})
53151
self.fn: Final[Any] = fn
54152
self.code: Final[str] = file.read_text()
55153
self.region_dialect: Final[str] = "cuda"
56154
self.lowered_region_dialect: Final[str] = "llvm"
57155
self.arg_dialect: Final[str] = "llvm"
58156
self.source_file: Final[str] = str(file)
157+
self.compiler = kwargs.get("compiler", None)
158+
self.target = kwargs.get("target", None)
159+
self.extern_file = kwargs.get("extern_file", None)
59160
self.extern_func_name = kwargs.get("extern_func_name", None)
60161
self.deferred: Final[bool] = kwargs.get("deferred", False)
61162
self.__triton_builtin__: Final[bool] = True
62163

164+
if self.compiler.lower() == "nvcc" and knobs.runtime.add_stages_inspection_hook is None:
165+
nvcc_cuda_hook = partial(make_cubin_inspection_hook, self)
166+
knobs.runtime.add_stages_inspection_hook = nvcc_cuda_hook
167+
63168
def register_pending_source(self, *, hint: str = "") -> str:
64169
if not self.extern_func_name:
65170
raise RuntimeError("deferred tle_raw CUDA source requires extern_func_name= "
@@ -116,6 +221,38 @@ def make_llvm(self, mlir_context) -> str:
116221
module = parse_llvm_ir(_sanitize_clang_ir(build.stdout.decode()), llvm_context, mlir_context)
117222
return f"{module}"
118223

224+
def make_bc(self, public_api_names=None):
225+
fbc_cuda_unopti = Path(self.source_file).with_suffix('.bc.unopti')
226+
fbc_cuda = Path(self.source_file).with_suffix('.bc')
227+
228+
build = subprocess.run([
229+
CLANG, "-c", "-x", "cuda", "--cuda-device-only",
230+
_get_cuda_gpu_arch(), "-emit-llvm",
231+
# "-O1",
232+
"-fcuda-flush-denormals-to-zero", *CLANG_FLAGS, "-o", fbc_cuda_unopti, self.source_file
233+
], capture_output=True)
234+
assert build.returncode == 0, (f"clang failed\nstderr:\n{build.stderr.decode()}")
235+
236+
if public_api_names is None:
237+
public_api_names = [self.extern_func_name]
238+
elif isinstance(public_api_names, str):
239+
public_api_names = [public_api_names]
240+
else:
241+
public_api_names = list(public_api_names)
242+
if not public_api_names or any(not name for name in public_api_names):
243+
raise ValueError("make_bc requires at least one public API name")
244+
public_api_list = ",".join(dict.fromkeys(public_api_names))
245+
246+
opt = subprocess.run([
247+
OPT, "--passes=internalize,inline,globaldce", f"-internalize-public-api-list={public_api_list}", "-o",
248+
fbc_cuda, fbc_cuda_unopti
249+
], capture_output=True)
250+
assert opt.returncode == 0, (f"opt failed\nstderr:\n{opt.stderr.decode()}")
251+
252+
if os.path.exists(fbc_cuda_unopti):
253+
os.remove(fbc_cuda_unopti)
254+
return fbc_cuda
255+
119256

120257
def compile_deferred_pending_source(entry: dict, *, context) -> str:
121258
source_text = entry["source"]

0 commit comments

Comments
 (0)