diff --git a/CONTRIBUTORS.md b/CONTRIBUTORS.md index c57f9a797..8149b1157 100644 --- a/CONTRIBUTORS.md +++ b/CONTRIBUTORS.md @@ -169,3 +169,4 @@ From oldest to newest contributor, we would like to thank: - [natinusala](https://github.com/natinusala) - [LJ](https://github.com/elle-j) - [Frank Leon Rose](https://github.com/frankleonrose) +- [Oguz Ulgen](https://github.com/oulgen) diff --git a/etc/config/helion.defaults.properties b/etc/config/helion.defaults.properties new file mode 100644 index 000000000..870e72f4c --- /dev/null +++ b/etc/config/helion.defaults.properties @@ -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 diff --git a/etc/scripts/helion_wrapper.py b/etc/scripts/helion_wrapper.py new file mode 100644 index 000000000..fe56f0388 --- /dev/null +++ b/etc/scripts/helion_wrapper.py @@ -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() + + diff --git a/examples/helion/default.py b/examples/helion/default.py new file mode 100644 index 000000000..d9abe41a5 --- /dev/null +++ b/examples/helion/default.py @@ -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) diff --git a/lib/compilers/_all.ts b/lib/compilers/_all.ts index d3e04694e..0fc20b218 100644 --- a/lib/compilers/_all.ts +++ b/lib/compilers/_all.ts @@ -92,6 +92,7 @@ export {GM2Compiler} from './gm2.js'; export {GnuCobolCompiler} from './gnucobol.js'; export {GolangCompiler} from './golang.js'; export {HaskellCompiler} from './haskell.js'; +export {HelionCompiler} from './helion.js'; export {HLSLCompiler} from './hlsl.js'; export {HookCompiler} from './hook.js'; export {HyloCompiler} from './hylo.js'; diff --git a/lib/compilers/helion.ts b/lib/compilers/helion.ts new file mode 100644 index 000000000..5990b821d --- /dev/null +++ b/lib/compilers/helion.ts @@ -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 --inputfile + return ['-I', this.compilerWrapperPath, '--outputfile', outputFilename, '--inputfile']; + } + + override getArgumentParserClass() { + return BaseParser; + } +} diff --git a/lib/languages.ts b/lib/languages.ts index 79c3602f0..71d69e157 100644 --- a/lib/languages.ts +++ b/lib/languages.ts @@ -488,6 +488,18 @@ const definitions: Record = { previewFilter: 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: { name: 'Hook', monaco: 'hook', diff --git a/public/logos/helion.png b/public/logos/helion.png new file mode 100644 index 000000000..35738ec32 Binary files /dev/null and b/public/logos/helion.png differ diff --git a/types/languages.interfaces.ts b/types/languages.interfaces.ts index 5b5ac9eb5..e6ff819b3 100644 --- a/types/languages.interfaces.ts +++ b/types/languages.interfaces.ts @@ -62,6 +62,7 @@ export type LanguageKey = | 'go' | 'haskell' | 'hlsl' + | 'helion' | 'hook' | 'hylo' | 'ispc'