mirror of
https://github.com/compiler-explorer/compiler-explorer.git
synced 2025-12-27 10:33:59 -05:00
Add Helion language support (#8206)
Helion is a Python-embedded domain-specific language (DSL) for authoring machine learning kernels, designed to compile down to Triton. https://github.com/pytorch/helion I mostly followed the Triton example, please let me know if there's more I need to do or anything I missed. infra PR: https://github.com/compiler-explorer/infra/pull/1875
This commit is contained in:
@@ -169,3 +169,4 @@ From oldest to newest contributor, we would like to thank:
|
|||||||
- [natinusala](https://github.com/natinusala)
|
- [natinusala](https://github.com/natinusala)
|
||||||
- [LJ](https://github.com/elle-j)
|
- [LJ](https://github.com/elle-j)
|
||||||
- [Frank Leon Rose](https://github.com/frankleonrose)
|
- [Frank Leon Rose](https://github.com/frankleonrose)
|
||||||
|
- [Oguz Ulgen](https://github.com/oulgen)
|
||||||
|
|||||||
15
etc/config/helion.defaults.properties
Normal file
15
etc/config/helion.defaults.properties
Normal file
@@ -0,0 +1,15 @@
|
|||||||
|
compilers=&helion
|
||||||
|
defaultCompiler=helion_020
|
||||||
|
compilerType=helion
|
||||||
|
interpreted=true
|
||||||
|
supportsBinary=false
|
||||||
|
supportsExecute=false
|
||||||
|
isSemVer=true
|
||||||
|
notification=Experimental Helion support on Compiler Explorer. Helion is a Python-based GPU kernel compiler that generates Triton code.
|
||||||
|
|
||||||
|
group.helion.compilers=helion_020
|
||||||
|
group.helion.groupName=Helion
|
||||||
|
|
||||||
|
compiler.helion_020.name=Helion 0.2.0
|
||||||
|
compiler.helion_020.exe=/opt/compiler-explorer/helion/v0.2.0/bin/python3
|
||||||
|
compiler.helion_020.semver=0.2.0
|
||||||
94
etc/scripts/helion_wrapper.py
Normal file
94
etc/scripts/helion_wrapper.py
Normal file
@@ -0,0 +1,94 @@
|
|||||||
|
# Copyright (c) 2025, Compiler Explorer Authors
|
||||||
|
# All rights reserved.
|
||||||
|
#
|
||||||
|
# Redistribution and use in source and binary forms, with or without
|
||||||
|
# modification, are permitted provided that the following conditions are met:
|
||||||
|
#
|
||||||
|
# * Redistributions of source code must retain the above copyright notice,
|
||||||
|
# this list of conditions and the following disclaimer.
|
||||||
|
# * Redistributions in binary form must reproduce the above copyright
|
||||||
|
# notice, this list of conditions and the following disclaimer in the
|
||||||
|
# documentation and/or other materials provided with the distribution.
|
||||||
|
#
|
||||||
|
# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
|
||||||
|
# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
|
||||||
|
# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
|
||||||
|
# ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
|
||||||
|
# LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
|
||||||
|
# CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
|
||||||
|
# SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
|
||||||
|
# INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
|
||||||
|
# CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
|
||||||
|
# ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
|
||||||
|
# POSSIBILITY OF SUCH DAMAGE.
|
||||||
|
|
||||||
|
from __future__ import annotations
|
||||||
|
|
||||||
|
import argparse
|
||||||
|
import contextlib
|
||||||
|
import importlib.util
|
||||||
|
import sys
|
||||||
|
|
||||||
|
|
||||||
|
def main() -> None:
|
||||||
|
parser = argparse.ArgumentParser(description="Output Triton code from public Helion kernels.")
|
||||||
|
parser.add_argument("--inputfile", required=True)
|
||||||
|
parser.add_argument("--outputfile", required=True)
|
||||||
|
args = parser.parse_args()
|
||||||
|
|
||||||
|
try:
|
||||||
|
import helion
|
||||||
|
from helion.runtime.kernel import Kernel
|
||||||
|
|
||||||
|
compiled_kernels: list[tuple[Kernel, object, str]] = []
|
||||||
|
|
||||||
|
# Patch kernel decorator to set autotune_effort='none' by default
|
||||||
|
original_kernel = helion.kernel
|
||||||
|
|
||||||
|
def patched_kernel(*args, **kwargs):
|
||||||
|
if 'config' not in kwargs and 'autotune_effort' not in kwargs:
|
||||||
|
kwargs['autotune_effort'] = 'none'
|
||||||
|
return original_kernel(*args, **kwargs)
|
||||||
|
|
||||||
|
helion.kernel = patched_kernel
|
||||||
|
|
||||||
|
original_call = Kernel.__call__
|
||||||
|
|
||||||
|
def patched_call(self, *call_args, **call_kwargs):
|
||||||
|
result = original_call(self, *call_args, **call_kwargs)
|
||||||
|
|
||||||
|
try:
|
||||||
|
bound = self.bind(call_args)
|
||||||
|
cfg = bound.config_spec.default_config()
|
||||||
|
triton_code = bound.to_triton_code(cfg)
|
||||||
|
compiled_kernels.append((self, call_args, triton_code))
|
||||||
|
except Exception:
|
||||||
|
pass
|
||||||
|
|
||||||
|
return result
|
||||||
|
|
||||||
|
Kernel.__call__ = patched_call
|
||||||
|
|
||||||
|
spec = importlib.util.spec_from_file_location("example", args.inputfile)
|
||||||
|
assert spec is not None and spec.loader is not None
|
||||||
|
mod = importlib.util.module_from_spec(spec)
|
||||||
|
spec.loader.exec_module(mod)
|
||||||
|
|
||||||
|
Kernel.__call__ = original_call
|
||||||
|
|
||||||
|
with open(args.outputfile, "w", encoding="utf-8") as out:
|
||||||
|
for kernel, args_used, triton_code in compiled_kernels:
|
||||||
|
out.write(triton_code)
|
||||||
|
out.write("\n\n")
|
||||||
|
|
||||||
|
except Exception as error:
|
||||||
|
messages = [m for m in (getattr(error, "args", None) or [str(error)])]
|
||||||
|
with contextlib.suppress(Exception):
|
||||||
|
sys.stderr.writelines([str(m) + "\n" for m in messages])
|
||||||
|
sys.exit(255)
|
||||||
|
|
||||||
|
|
||||||
|
if __name__ == "__main__":
|
||||||
|
main()
|
||||||
|
|
||||||
|
|
||||||
29
examples/helion/default.py
Normal file
29
examples/helion/default.py
Normal file
@@ -0,0 +1,29 @@
|
|||||||
|
import torch
|
||||||
|
import helion
|
||||||
|
import helion.language as hl
|
||||||
|
|
||||||
|
|
||||||
|
@helion.kernel(config=helion.Config(
|
||||||
|
block_sizes=[],
|
||||||
|
indexing='pointer',
|
||||||
|
load_eviction_policies=['', ''],
|
||||||
|
num_stages=2,
|
||||||
|
num_warps=4,
|
||||||
|
pid_type='flat',
|
||||||
|
range_flattens=[None],
|
||||||
|
range_multi_buffers=[None],
|
||||||
|
range_num_stages=[0],
|
||||||
|
range_unroll_factors=[0],
|
||||||
|
range_warp_specializes=[]
|
||||||
|
), static_shapes=True)
|
||||||
|
def add(x: torch.Tensor, y: torch.Tensor, scale: float) -> torch.Tensor:
|
||||||
|
out = torch.empty_like(x)
|
||||||
|
for idx in hl.grid(x.size()):
|
||||||
|
out[idx] = (x[idx] + y[idx]) * scale
|
||||||
|
return out
|
||||||
|
|
||||||
|
|
||||||
|
x = torch.rand(1024).cuda()
|
||||||
|
y = torch.rand(1024).cuda()
|
||||||
|
scale = 2.0
|
||||||
|
result = add(x, y, scale)
|
||||||
@@ -92,6 +92,7 @@ export {GM2Compiler} from './gm2.js';
|
|||||||
export {GnuCobolCompiler} from './gnucobol.js';
|
export {GnuCobolCompiler} from './gnucobol.js';
|
||||||
export {GolangCompiler} from './golang.js';
|
export {GolangCompiler} from './golang.js';
|
||||||
export {HaskellCompiler} from './haskell.js';
|
export {HaskellCompiler} from './haskell.js';
|
||||||
|
export {HelionCompiler} from './helion.js';
|
||||||
export {HLSLCompiler} from './hlsl.js';
|
export {HLSLCompiler} from './hlsl.js';
|
||||||
export {HookCompiler} from './hook.js';
|
export {HookCompiler} from './hook.js';
|
||||||
export {HyloCompiler} from './hylo.js';
|
export {HyloCompiler} from './hylo.js';
|
||||||
|
|||||||
55
lib/compilers/helion.ts
Normal file
55
lib/compilers/helion.ts
Normal file
@@ -0,0 +1,55 @@
|
|||||||
|
// Copyright (c) 2025, Compiler Explorer Authors
|
||||||
|
// All rights reserved.
|
||||||
|
//
|
||||||
|
// Redistribution and use in source and binary forms, with or without
|
||||||
|
// modification, are permitted provided that the following conditions are met:
|
||||||
|
//
|
||||||
|
// * Redistributions of source code must retain the above copyright notice,
|
||||||
|
// this list of conditions and the following disclaimer.
|
||||||
|
// * Redistributions in binary form must reproduce the above copyright
|
||||||
|
// notice, this list of conditions and the following disclaimer in the
|
||||||
|
// documentation and/or other materials provided with the distribution.
|
||||||
|
//
|
||||||
|
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
|
||||||
|
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
|
||||||
|
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
|
||||||
|
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
|
||||||
|
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
|
||||||
|
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
|
||||||
|
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
|
||||||
|
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
|
||||||
|
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
|
||||||
|
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
|
||||||
|
// POSSIBILITY OF SUCH DAMAGE.
|
||||||
|
|
||||||
|
import type {PreliminaryCompilerInfo} from '../../types/compiler.interfaces.js';
|
||||||
|
import type {ParseFiltersAndOutputOptions} from '../../types/features/filters.interfaces.js';
|
||||||
|
import {BaseCompiler} from '../base-compiler.js';
|
||||||
|
import {CompilationEnvironment} from '../compilation-env.js';
|
||||||
|
import {resolvePathFromAppRoot} from '../utils.js';
|
||||||
|
|
||||||
|
import {BaseParser} from './argument-parsers.js';
|
||||||
|
|
||||||
|
export class HelionCompiler extends BaseCompiler {
|
||||||
|
private compilerWrapperPath: string;
|
||||||
|
|
||||||
|
static get key() {
|
||||||
|
return 'helion';
|
||||||
|
}
|
||||||
|
|
||||||
|
constructor(compilerInfo: PreliminaryCompilerInfo, env: CompilationEnvironment) {
|
||||||
|
super(compilerInfo, env);
|
||||||
|
this.compilerWrapperPath =
|
||||||
|
this.compilerProps('compilerWrapper', '') || resolvePathFromAppRoot('etc', 'scripts', 'helion_wrapper.py');
|
||||||
|
// We are producing Triton code as the primary output; use default asm parser handling
|
||||||
|
}
|
||||||
|
|
||||||
|
override optionsForFilter(filters: ParseFiltersAndOutputOptions, outputFilename: string): string[] {
|
||||||
|
// Wrapper contract: python -I helion_wrapper.py --outputfile <file> --inputfile <src>
|
||||||
|
return ['-I', this.compilerWrapperPath, '--outputfile', outputFilename, '--inputfile'];
|
||||||
|
}
|
||||||
|
|
||||||
|
override getArgumentParserClass() {
|
||||||
|
return BaseParser;
|
||||||
|
}
|
||||||
|
}
|
||||||
@@ -488,6 +488,18 @@ const definitions: Record<LanguageKey, LanguageDefinition> = {
|
|||||||
previewFilter: null,
|
previewFilter: null,
|
||||||
monacoDisassembly: null,
|
monacoDisassembly: null,
|
||||||
},
|
},
|
||||||
|
helion: {
|
||||||
|
name: 'Helion',
|
||||||
|
monaco: 'python',
|
||||||
|
extensions: ['.py'],
|
||||||
|
alias: [],
|
||||||
|
logoFilename: 'helion.png',
|
||||||
|
logoFilenameDark: 'helion.png',
|
||||||
|
formatter: null,
|
||||||
|
previewFilter: null,
|
||||||
|
monacoDisassembly: null,
|
||||||
|
digitSeparator: '_',
|
||||||
|
},
|
||||||
hook: {
|
hook: {
|
||||||
name: 'Hook',
|
name: 'Hook',
|
||||||
monaco: 'hook',
|
monaco: 'hook',
|
||||||
|
|||||||
BIN
public/logos/helion.png
Normal file
BIN
public/logos/helion.png
Normal file
Binary file not shown.
|
After Width: | Height: | Size: 55 KiB |
@@ -62,6 +62,7 @@ export type LanguageKey =
|
|||||||
| 'go'
|
| 'go'
|
||||||
| 'haskell'
|
| 'haskell'
|
||||||
| 'hlsl'
|
| 'hlsl'
|
||||||
|
| 'helion'
|
||||||
| 'hook'
|
| 'hook'
|
||||||
| 'hylo'
|
| 'hylo'
|
||||||
| 'ispc'
|
| 'ispc'
|
||||||
|
|||||||
Reference in New Issue
Block a user