Add files using upload-large-folder tool

.local/share/Trash/info/train_002.bin.trashinfo

@@ -0,0 +1,3 @@
+[Trash Info]
+Path=/root/data/fineweb/train_002.bin
+DeletionDate=2024-09-26T05:50:34

.local/share/Trash/info/train_004.bin.trashinfo

@@ -0,0 +1,3 @@
+[Trash Info]
+Path=/root/data/fineweb/train_004.bin
+DeletionDate=2024-09-26T05:50:34

.local/share/jupyter/nbextensions/ruler/main.js

@@ -0,0 +1,124 @@
+// Add rulers to codecells
+define([
+    'base/js/namespace',
+    'base/js/events',
+    'services/config',
+    'notebook/js/codecell',
+    'codemirror/lib/codemirror',
+    'codemirror/addon/display/rulers'
+], function (Jupyter, events, configmod, codecell, codemirror) {
+    "use strict";
+
+    var log_prefix = '[ruler]';
+
+    // define default config parameter values
+    var params = {
+        ruler_column: [78],
+        ruler_color: ["#ff0000"],
+        ruler_linestyle: ["dashed"],
+        ruler_do_css_patch: false
+    };
+
+
+    var rulers = [];
+
+    var isNumber = function (n) {
+        return !isNaN(parseFloat(n)) && isFinite(n);
+    };
+
+    // updates default params with any specified in the provided config data
+    var update_params = function (config_data) {
+        for (var key in params) {
+            if (config_data.hasOwnProperty(key)) {
+                params[key] = config_data[key];
+            }
+        }
+    };
+
+    var on_config_loaded = function () {
+
+        if (Jupyter.notebook !== undefined) {
+            var i, config = Jupyter.notebook.config;
+        } else {
+            var i, config = Jupyter.editor.config;
+        }
+
+        if (config.data.hasOwnProperty('ruler_color') && config.data.ruler_color.length > 0) {
+            params.ruler_color = config.data.ruler_color;
+        }
+
+        if (config.data.hasOwnProperty('ruler_column')) {
+            var new_columns = [];
+            for (i in config.data.ruler_column) {
+                if (isNumber(config.data.ruler_column[i])) {
+                    new_columns.push(config.data.ruler_column[i]);
+                }
+            }
+            if (new_columns.length > 0) {
+                params.ruler_column = new_columns;
+            }
+        }
+
+        if (config.data.hasOwnProperty('ruler_linestyle') && config.data.ruler_linestyle.length > 0) {
+            params.ruler_linestyle = config.data.ruler_linestyle;
+        }
+
+        for (i in params.ruler_column) {
+            rulers.push({
+                color: params.ruler_color[i % params.ruler_color.length],
+                column: params.ruler_column[i],
+                lineStyle: params.ruler_linestyle[i % params.ruler_linestyle.length]
+            });
+        }
+        console.debug(log_prefix, 'ruler specs:', rulers);
+
+        if (Jupyter.notebook !== undefined) {
+            var i, config = Jupyter.notebook.config;
+
+            // Change default for new cells
+            codecell.CodeCell.options_default.cm_config.rulers = rulers;
+            // Apply to any already-existing cells
+            var cells = Jupyter.notebook.get_cells().forEach(function (cell) {
+                if (cell instanceof codecell.CodeCell) {
+                    cell.code_mirror.setOption('rulers', rulers);
+                }
+            });
+
+        }
+        else {
+            Jupyter.editor.codemirror.setOption('rulers', rulers);
+        }
+    };
+
+    var load_extension = function () {
+
+        // first, check which view we're in, in order to decide whether to load
+        var conf_sect;
+        if (Jupyter.notebook) {
+            // we're in notebook view
+            conf_sect = Jupyter.notebook.config;
+        }
+        else if (Jupyter.editor) {
+            // we're in file-editor view
+            conf_sect = Jupyter.editor.config;
+        }
+        else {
+            // we're some other view like dashboard, terminal, etc, so bail now
+            return;
+        }
+
+        conf_sect.loaded
+            .then(function () {
+                update_params(conf_sect.data);
+            })
+            .then(on_config_loaded)
+            .catch(function on_error(reason) {
+                console.warn(log_prefix, 'error:', reason);
+            });
+    };
+
+    var extension = {
+        load_ipython_extension: load_extension
+    };
+    return extension;
+});

.local/share/jupyter/nbextensions/runtools/readme.md

@@ -0,0 +1,119 @@
+Runtools
+========
+Runtools provide a number of additional functions for working with code cells in the IPython notebook:
+
+Code Cell Execution
+-------------------
+
+* Execute a single cell
+* Execute from top cell to currently selected cell
+* Execute from currently selected cell to bottom cell
+* Execute all cells
+* Execute all cells, ignore exceptions (requires [ipython/pull/6521](https://github.com/ipython/ipython/pull/6521))
+* Execute marked code cells (cells with green gutter area are marked)
+* Stop execution (duplicate to standard toolbar button)
+
+When executing marked cells, they are put in a execution list, and
+executed in order. The execution list can be modified by unmarking
+a cell not yet run. The execution list can be stopped by clicking on
+`stop execution`. Execution of the currently running cell can be stopped
+by pressing `stop execution` twice.
+
+Code Cell Marking
+-----------------
+
+* Mark one or more code cell
+
+
+Code Cell Display
+-----------------
+
+* Hide or show input (i.e. the source code) of marked code cells
+* Hide or show output of marked code cells
+
+
+Description
+-----------
+
+The *runtools* extension adds a button to turn on/off a floating toolbar:   
+![](icon.png)
+
+This adds Code execution buttons:   
+![](runtools_execute.png)
+
+Codecells can be marked by clicking on the gutter of a codecell or by clicking on the markers toolbar:   
+![](runtools_marker.png)
+
+Marked codecells can be locked to read-only mode and moved upd and down:   
+![](runtools_lock.png)
+
+The input and output areas of marked codecells can be hidden:   
+![](runtools_show_hide.png)
+
+A IPython notebook with marked cells looks like this:
+![](runtools_nb.png)
+
+
+Demo
+----
+
+![](demo.gif)
+
+
+Internals
+---------
+
+New metadata elements added to each cell:
+* `cell.metadata.hide_input` - hide input field of the cell
+* `cell.metadata.hide_output` - hide output field of the cell
+* `cell.metadata.run_control.marked` - mark a codecell
+
+To export a notebook with hidden input/output fields, the custom template `hide_input_output.tpl` is required.
+It should have been installed in the `templates` folder.
+You can find the `templates` folder of `jupyter_contrib_nbextensions` from python using
+
+```python
+from jupyter_contrib_nbextensions.nbconvert_support import templates_directory
+print(templates_directory())
+```
+
+The template needs to be in a path where nbconvert can find it. This can be your local path or specified in 
+`jupyter_nbconvert_config` or `jupyter_notebook_config` as `c.Exporter.extra_template_paths`, see [Jupyter docs](https://jupyter-notebook.readthedocs.io/en/latest/config.html).
+
+For HTML export a template is provided as `nbextensions.tpl` in the `jupyter_contrib_nbextensions` templates directory. Alternatively you can create your own template:
+```
+{%- extends 'full.tpl' -%}
+
+{% block input_group -%}
+{%- if cell.metadata.hide_input -%}
+{%- else -%}
+{{ super() }}
+{%- endif -%}
+{% endblock input_group %}
+
+{% block output_group -%}
+{%- if cell.metadata.hide_output -%}
+{%- else -%}
+{{ super() }}
+{%- endif -%}
+{% endblock output_group %}
+```
+
+For LaTeX export a different template is required, which is included as `nbextensions.tplx` in the `jupyter_contrib_nbextensions` templates directory. Alternatively you can create your own template:
+```
+((- extends 'report.tplx' -))
+
+((* block input_group -))
+((- if cell.metadata.hide_input -))
+((- else -))
+((( super() )))
+((- endif -))
+(( endblock input_group *))
+
+((* block output_group -))
+((- if cell.metadata.hide_output -))
+((- else -))
+((( super() )))
+((- endif -))
+(( endblock output_group *))
+```

.local/share/jupyter/nbextensions/scratchpad/scratchpad.yaml

@@ -0,0 +1,6 @@
+Type: Jupyter Notebook Extension
+Name: Scratchpad
+Description: Adds a scratchpad cell to Jupyter notebook.
+Link: README.md
+Main: main.js
+Compatibility: 4.x, 5.x

.local/share/jupyter/nbextensions/select_keymap/README.md

@@ -0,0 +1,14 @@
+Select CodeMirror Keymap
+=======
+
+This extension lets you choose between the available CodeMirror keymaps: default, emacs, vim, and sublime.
+
+There's a chance that this might cause key conflicts, especially with other extensions.
+
+Most browsers consume some of the global keybindings like `Ctrl+n`. [The Menu Wizard add-on for Firefox](https://addons.mozilla.org/en-US/firefox/addon/s3menu-wizard/) allows you to disable some of the global key shortcuts, thus passing the keys through to CodeMirror.
+
+![Demo](select_keymap.png)
+
+Based on:
+ * [jupyter-emacskeys](https://github.com/rmcgibbo/jupyter-emacskeys)
+ * [notebook_input_mode](https://github.com/asford/notebook_input_mode)

.local/share/jupyter/nbextensions/skill/main.js

@@ -0,0 +1,14 @@
+define(function() {
+	"use strict";
+	// jupyter nbextensions must export a load_ipython_extension function to
+	// avoid throwing an error. Also, loading the module should do nothing
+	// unless the function is called, so we wrap requiring the codemirror mode
+	// in the load call.
+	return {
+		load_ipython_extension: function () {
+			requirejs(['./skill'], function () {
+				console.log('[SKILL Syntax] loaded');
+			});
+		}
+	};
+});

.local/share/jupyter/nbextensions/skill/skill.yaml

@@ -0,0 +1,6 @@
+Type: IPython Notebook Extension
+Name: SKILL Syntax
+Description: Enable SKILL syntax support for CodeMirror
+Link: README.md
+Main: main.js
+Compatibility: 4.x, 5.x

.local/share/jupyter/nbextensions/splitcell/splitcell.js

@@ -0,0 +1,101 @@
+// Allow for split cells in jupyter notebooks
+
+define([
+	'base/js/namespace',
+		'base/js/events'
+], function (
+	Jupyter,
+	events
+) {
+	"use strict";
+
+	//define default config parameter values
+	var params = {
+		toggle_cell_style_keybinding : 'shift-s'
+	};
+
+	//updates default params with any specified in the server's config
+	var update_params = function(){
+		var config = Jupyter.notebook.config;
+		for (var key in params){
+			if (config.data.hasOwnProperty(key)){
+				params[key] = config.data[key];
+			}
+		}
+	};
+
+	 var setup = function (){
+		// update defaults
+		update_params();
+
+		//register actions with ActionHandler instance
+		var prefix = 'auto';
+		var name = 'toggle-cell-style';
+		var action = {
+			icon : 'fa-arrows-h',
+			help : 'Toggle split/centered cell style',
+			help_index : 'eb',
+			id : 'split_cells',
+			handler : toggle_cell_style
+		};
+
+		var action_full_name = Jupyter.keyboard_manager.actions.register(action, name, prefix);
+
+		//define keyboard shortucts
+		var command_mode_shortcuts = {};
+		command_mode_shortcuts[params.toggle_cell_style_keybinding] =  action_full_name;
+
+		//register keyboard shortucts with keyboard_manager
+		Jupyter.notebook.keyboard_manager.command_shortcuts.add_shortcuts(command_mode_shortcuts);
+		Jupyter.toolbar.add_buttons_group([action_full_name]);
+	};
+
+
+	var toggle_cell_style = function(){
+		var cell = Jupyter.notebook.get_selected_cell();
+		if (!("cell_style" in cell.metadata)){cell.metadata.cell_style = 'split';}
+		else if (cell.metadata.cell_style == 'center'){cell.metadata.cell_style = 'split';}
+		else {cell.metadata.cell_style = 'center';}
+
+		update_cell_style_element(cell);
+	};
+
+	var get_cell_style_html = function(cell_style){
+		console.log(cell_style);
+		if (cell_style == "split") 
+			{return "float:left; width:50%;";}
+		return "width:100%;";
+	};
+
+	var update_cell_style_element = function(cell){
+		var cell_style_html = get_cell_style_html(cell.metadata.cell_style);
+		cell.element.attr('style', cell_style_html);
+	};
+
+	function initialize () {
+		// On Load lets set the cell styles correctly
+		var cells = Jupyter.notebook.get_cells();
+		var ncells = Jupyter.notebook.ncells();
+
+		for (var i=0; i<ncells; i++){
+			var cell = cells[i];
+			if ("cell_style" in cell.metadata){
+				update_cell_style_element(cell, cell.metadata.cell_style);
+			}
+		 }
+	}
+
+	var load_extension = function() {
+		Jupyter.notebook.config.loaded.then(setup);
+
+		if (Jupyter.notebook !== undefined && Jupyter.notebook._fully_loaded) {
+			// notebook already loaded. Update directly
+			initialize();
+		}
+		events.on("notebook_loaded.Notebook", initialize);
+	};
+
+	return {
+		load_ipython_extension : load_extension
+	};
+});

.local/share/jupyter/nbextensions/varInspector/README.md

@@ -0,0 +1,36 @@
+# Variable Inspector
+
+## Description and main features
+
+The Variable Inspector extension, which currently supports python and R kernels,  enables to collect all defined variables and display them in a floating window. The window not only display the name of variables but also  their type, size in memory and content. The columns are sortable. The window is draggable, resizable, collapsable. The list of displayed variables is automatically updated at each cell execution. Variables can be deleted from workspace by clicking a link. Position and state (displayed/collapsed) are stored in the notebook's metadata and restored at startup.
+
+The extension supports multiple kernels. To add support for a new kernel, one has to
+- provide a library which loads required modules and define a function which lists all variables, together with their name, type, size and content. The output of this function must be a JSON representation of a list of objects (one for each variable) with keys 'varName','varType', 'varSize', 'varContent',
+- provide the command for deleting a variable, as `delete_cmd_prefix` and `delete_cmd_postfix`, eg. for `rm(variable)`, specify `rm(` and `)`.
+- give the command to refresh the list of variables (usually this is a call to the function defined in the library above). This information can be provided either in the source file or in the yaml config file.
+
+In any case, contributions to support further kernels will be very welcome!
+
+#### Demo:
+![](demo.gif)
+
+
+## Configuration
+The initial configuration can be given using the IPython-contrib nbextensions facility. It includes:
+
+- varInspector.window_display - Display at startup or not (default: false)
+- varInspector.cols.lenName: (and .lenType, .lenVar) - Width of columns (actually the max number of character to display in each column)
+- varInspector.kernels_config - json object defining the kernels specific code and commands.
+
+
+## Notes
+- The displayed size of variables use the `getsizeof()` python method. This method doesn't work for all types, so the reported size is to be considered with some caution. The extension includes some code to correctly return the size of numpy arrays, pandas Series and DataFrame but the size for some other types may be incorrect.
+- The extension builds on some code provided [here](https://github.com/jupyter-widgets/ipywidgets/blob/master/docs/source/examples/Variable%20Inspector.ipynb)  (essentially the `_fill` method)
+- The extension uses Christian Bach's [table sorter jquery plugin](https://github.com/christianbach/tablesorter). License file is included.
+
+
+## History
+
+- @jfbercher march 22, 2017 -- initial release
+- @jfbercher april 03, 2017 -- multiple kernel support. added support for R kernels.
+- @jfbercher june 30, 2017 -- fixed #1014 (use of `%reset` with IPython kernel) and #1015 printing with python 2 kernel.

.local/share/jupyter/nbextensions/varInspector/tablesorter_LICENSE.txt

@@ -0,0 +1,21 @@
+The MIT License (MIT)
+
+Copyright (c) 2014 Christian Bach
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all
+copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+SOFTWARE.

.local/share/jupyter/nbextensions/varInspector/var_list.py

@@ -0,0 +1,63 @@
+import json
+from sys import getsizeof
+
+from IPython import get_ipython
+from IPython.core.magics.namespace import NamespaceMagics
+_nms = NamespaceMagics()
+_Jupyter = get_ipython()
+_nms.shell = _Jupyter.kernel.shell
+
+try:
+    import numpy as np
+except ImportError:
+    pass    
+
+def _getsizeof(x):
+    # return the size of variable x. Amended version of sys.getsizeof
+    # which also supports ndarray, Series and DataFrame
+    if type(x).__name__ in ['ndarray', 'Series']:
+        return x.nbytes
+    elif type(x).__name__ == 'DataFrame':
+        return x.memory_usage().sum()
+    else:
+        return getsizeof(x)
+
+def _getshapeof(x):
+    #returns the shape of x if it has one
+    #returns None otherwise - might want to return an empty string for an empty column
+    try:
+        return x.shape
+    except AttributeError: #x does not have a shape
+        return None
+
+def _getcontentof(x):
+    length = 150
+    if type(x).__name__ == 'DataFrame':
+        colnames = ', '.join(x.columns.map(str))
+        content = "Column names: %s" % colnames
+    elif type(x).__name__ == 'Series':
+        content = "Series [%d rows]" % x.shape
+    elif type(x).__name__ == 'ndarray':
+        content = x.__repr__()
+    else:
+        if hasattr(x, '__len__'):
+            if len(x) > length:
+                content = str(x[:length])
+        else:
+            content = str(x)
+        if len(content) > 150:
+            return content[:150] + " ..."
+    return content
+
+def var_dic_list():
+    types_to_exclude = ['module', 'function', 'builtin_function_or_method',
+                        'instance', '_Feature', 'type', 'ufunc']
+    values = _nms.who_ls()
+    vardic = [{'varName': v, 'varType': type(eval(v)).__name__, 'varSize': str(_getsizeof(eval(v))), 'varShape': str(_getshapeof(eval(v))) if _getshapeof(eval(v)) else '', 'varContent': _getcontentof(eval(v)) }  # noqa
+    
+    for v in values if (v not in ['_html', '_nms', 'NamespaceMagics', '_Jupyter']) & (type(eval(v)).__name__ not in types_to_exclude)] # noqa 
+    return json.dumps(vardic)
+
+
+# command to refresh the list of variables
+print(var_dic_list())

.local/share/jupyter/nbextensions/varInspector/var_list.r

@@ -0,0 +1,17 @@
+library(jsonlite)
+var_dic_list = function(){ 
+    ll = ls(.GlobalEnv, all.names = FALSE)
+    varList=list()
+    iter = 1
+        for (k in ll){
+            if (class(get(k))!='function'){
+            class = class(get(k)); rk = capture.output(str(get(k))); size =  object.size(get(k)); sk = substr(get(k),0, 200); 
+            #  [{'varName':v, 'varType': type(eval(v)).__name__, 'varSize': _getsizeof(eval(v)), 'varContent': str(eval(v))[:200]} 
+            l = list(varName = k, varType = class, varSize = size, varContent = sk)
+            varList[[iter]] = l
+            # print(l)
+            iter = iter + 1}
+        }
+return(toJSON(varList, simplifyVector = FALSE, force=TRUE))
+    }
+cat(var_dic_list())

.local/share/jupyter/nbextensions/zenmode/README.md

@@ -0,0 +1,4 @@
+Zenmode
+=======
+
+A little extension to give Zenmode functionality to the IPython notebook

.local/share/jupyter/nbextensions/zenmode/main.css

@@ -0,0 +1,34 @@
+.navbar-inner {
+    opacity: 0.5;
+    -webkit-transition: opacity 0.3s ease-in-out;
+    -moz-transition: opacity 0.3s ease-in-out;
+    -o-transition: opacity 0.3s ease-in-out;
+    transition: opacity 0.3s ease-in-out;
+}
+
+.navbar-inner:hover {
+    opacity: 1.0;
+}
+
+#maintoolbar .navbar-text {
+    display: none !important;
+}
+
+#notebook-container {
+    background-color: rgba(255, 255, 255, 0);
+}
+
+/*
+.cell {
+    background-color: rgb(255, 255, 255);
+}
+
+.CodeMirror {
+    background: #F8FCCF;
+}
+
+div.input_area {
+    margin: 2px;
+    border: none;
+}
+*/

.local/share/jupyter/runtime/jpserver-434.json

@@ -0,0 +1,13 @@
+{
+  "base_url": "/",
+  "hostname": "0.0.0.0",
+  "password": false,
+  "pid": 434,
+  "port": 8080,
+  "root_dir": "/root",
+  "secure": true,
+  "sock": "",
+  "token": "5a434251505375f2b42435914de608ef3450739f4e14b0be1cfeae3b7364239e",
+  "url": "https://184d1c0992ce:8080/",
+  "version": "2.12.5"
+}

.triton/dump/0f43b9f3b1f9407355f6ad39f7d56744/triton_.ttgir

@@ -0,0 +1,60 @@
+#blocked = #triton_gpu.blocked<{sizePerThread = [1, 1], threadsPerWarp = [4, 8], warpsPerCTA = [1, 4], order = [0, 1], CTAsPerCGA = [1, 1], CTASplitNum = [1, 1], CTAOrder = [1, 0]}>
+#blocked1 = #triton_gpu.blocked<{sizePerThread = [4, 1], threadsPerWarp = [1, 32], warpsPerCTA = [1, 4], order = [0, 1], CTAsPerCGA = [1, 1], CTASplitNum = [1, 1], CTAOrder = [1, 0]}>
+module attributes {"triton_gpu.compute-capability" = 89 : i32, "triton_gpu.num-ctas" = 1 : i32, "triton_gpu.num-warps" = 4 : i32, "triton_gpu.threads-per-warp" = 32 : i32} {
+  tt.func public @triton__0d1d2d3de4e(%arg0: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<i64, 1> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg3: i32 {tt.divisibility = 16 : i32, tt.max_divisibility = 16 : i32}, %arg4: i32 {tt.max_divisibility = 8 : i32}) attributes {noinline = false} {
+    %cst = arith.constant dense<256> : tensor<4x1xi64, #blocked>
+    %cst_0 = arith.constant dense<0> : tensor<4x1xi64, #blocked>
+    %cst_1 = arith.constant dense<512> : tensor<4x1xi64, #blocked>
+    %cst_2 = arith.constant dense<256> : tensor<4x1xi32, #blocked>
+    %cst_3 = arith.constant dense<131072> : tensor<1x128xi32, #blocked1>
+    %cst_4 = arith.constant dense<120> : tensor<1x128xi32, #blocked1>
+    %cst_5 = arith.constant dense<0.000000e+00> : tensor<4x128xf32, #blocked1>
+    %cst_6 = arith.constant dense<true> : tensor<4x1xi1, #blocked>
+    %c4_i32 = arith.constant 4 : i32
+    %0 = tt.get_program_id x : i32
+    %1 = arith.muli %0, %c4_i32 : i32
+    %2 = tt.make_range {end = 4 : i32, start = 0 : i32} : tensor<4xi32, #triton_gpu.slice<{dim = 1, parent = #blocked1}>>
+    %3 = tt.make_range {end = 4 : i32, start = 0 : i32} : tensor<4xi32, #triton_gpu.slice<{dim = 1, parent = #blocked}>>
+    %4 = tt.expand_dims %2 {axis = 1 : i32} : (tensor<4xi32, #triton_gpu.slice<{dim = 1, parent = #blocked1}>>) -> tensor<4x1xi32, #blocked1>
+    %5 = tt.expand_dims %3 {axis = 1 : i32} : (tensor<4xi32, #triton_gpu.slice<{dim = 1, parent = #blocked}>>) -> tensor<4x1xi32, #blocked>
+    %6 = tt.splat %1 : (i32) -> tensor<4x1xi32, #blocked1>
+    %7 = tt.splat %1 : (i32) -> tensor<4x1xi32, #blocked>
+    %8 = arith.addi %6, %4 : tensor<4x1xi32, #blocked1>
+    %9 = arith.addi %7, %5 : tensor<4x1xi32, #blocked>
+    %10 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32, #triton_gpu.slice<{dim = 0, parent = #blocked1}>>
+    %11 = tt.expand_dims %10 {axis = 0 : i32} : (tensor<128xi32, #triton_gpu.slice<{dim = 0, parent = #blocked1}>>) -> tensor<1x128xi32, #blocked1>
+    %12 = arith.cmpi slt, %11, %cst_4 : tensor<1x128xi32, #blocked1>
+    %13 = arith.muli %11, %cst_3 : tensor<1x128xi32, #blocked1>
+    %14 = tt.broadcast %8 : (tensor<4x1xi32, #blocked1>) -> tensor<4x128xi32, #blocked1>
+    %15 = tt.broadcast %13 : (tensor<1x128xi32, #blocked1>) -> tensor<4x128xi32, #blocked1>
+    %16 = arith.addi %14, %15 : tensor<4x128xi32, #blocked1>
+    %17 = tt.splat %arg0 : (!tt.ptr<f32, 1>) -> tensor<4x128x!tt.ptr<f32, 1>, #blocked1>
+    %18 = tt.addptr %17, %16 : tensor<4x128x!tt.ptr<f32, 1>, #blocked1>, tensor<4x128xi32, #blocked1>
+    %19 = tt.broadcast %12 : (tensor<1x128xi1, #blocked1>) -> tensor<4x128xi1, #blocked1>
+    %20 = tt.load %18, %19, %cst_5 {cache = 1 : i32, evict = 2 : i32, isVolatile = false} : tensor<4x128xf32, #blocked1>
+    %21 = arith.addf %20, %cst_5 : tensor<4x128xf32, #blocked1>
+    %22 = arith.select %19, %21, %cst_5 : tensor<4x128xi1, #blocked1>, tensor<4x128xf32, #blocked1>
+    %23 = "tt.reduce"(%22) <{axis = 1 : i32}> ({
+    ^bb0(%arg5: f32, %arg6: f32):
+      %40 = arith.addf %arg5, %arg6 : f32
+      tt.reduce.return %40 : f32
+    }) : (tensor<4x128xf32, #blocked1>) -> tensor<4xf32, #triton_gpu.slice<{dim = 1, parent = #blocked1}>>
+    %24 = triton_gpu.convert_layout %23 : (tensor<4xf32, #triton_gpu.slice<{dim = 1, parent = #blocked1}>>) -> tensor<4xf32, #triton_gpu.slice<{dim = 1, parent = #blocked}>>
+    %25 = tt.expand_dims %24 {axis = 1 : i32} : (tensor<4xf32, #triton_gpu.slice<{dim = 1, parent = #blocked}>>) -> tensor<4x1xf32, #blocked>
+    %26 = arith.divsi %9, %cst_2 : tensor<4x1xi32, #blocked>
+    %27 = arith.remsi %9, %cst_2 : tensor<4x1xi32, #blocked>
+    %28 = tt.splat %arg1 : (!tt.ptr<i64, 1>) -> tensor<4x1x!tt.ptr<i64, 1>, #blocked>
+    %29 = tt.addptr %28, %26 : tensor<4x1x!tt.ptr<i64, 1>, #blocked>, tensor<4x1xi32, #blocked>
+    %30 = tt.load %29 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<4x1xi64, #blocked>
+    %31 = arith.addi %30, %cst_1 : tensor<4x1xi64, #blocked>
+    %32 = arith.cmpi slt, %30, %cst_0 : tensor<4x1xi64, #blocked>
+    %33 = arith.select %32, %31, %30 : tensor<4x1xi1, #blocked>, tensor<4x1xi64, #blocked>
+    %34 = arith.muli %33, %cst : tensor<4x1xi64, #blocked>
+    %35 = arith.extsi %27 : tensor<4x1xi32, #blocked> to tensor<4x1xi64, #blocked>
+    %36 = arith.addi %35, %34 : tensor<4x1xi64, #blocked>
+    %37 = tt.splat %arg2 : (!tt.ptr<f32, 1>) -> tensor<4x1x!tt.ptr<f32, 1>, #blocked>
+    %38 = tt.addptr %37, %36 : tensor<4x1x!tt.ptr<f32, 1>, #blocked>, tensor<4x1xi64, #blocked>
+    %39 = "tt.atomic_rmw"(%38, %25, %cst_6) <{atomic_rmw_op = 5 : i32, scope = 1 : i32, sem = 4 : i32}> : (tensor<4x1x!tt.ptr<f32, 1>, #blocked>, tensor<4x1xf32, #blocked>, tensor<4x1xi1, #blocked>) -> tensor<4x1xf32, #blocked>
+    tt.return
+  }
+}

.triton/dump/199215289adb100508718a5a762ba4d7/triton_.ptx

@@ -0,0 +1,453 @@
+//
+// Generated by LLVM NVPTX Back-End
+//
+
+.version 8.2
+.target sm_89
+.address_size 64
+
+	// .globl	triton__0d1d2de
+.extern .func __assertfail
+(
+	.param .b64 __assertfail_param_0,
+	.param .b64 __assertfail_param_1,
+	.param .b32 __assertfail_param_2,
+	.param .b64 __assertfail_param_3,
+	.param .b64 __assertfail_param_4
+)
+;
+.global .align 1 .b8 assertFunc_0[25] = {95, 99, 97, 108, 108, 95, 119, 105, 116, 104, 95, 102, 114, 97, 109, 101, 115, 95, 114, 101, 109, 111, 118, 101, 100};
+.global .align 1 .b8 assertFile_0[38] = {60, 102, 114, 111, 122, 101, 110, 32, 105, 109, 112, 111, 114, 116, 108, 105, 98, 46, 95, 98, 111, 111, 116, 115, 116, 114, 97, 112, 95, 101, 120, 116, 101, 114, 110, 97, 108, 62};
+.global .align 1 .b8 assertMessage_0[38] = {105, 110, 100, 101, 120, 32, 111, 117, 116, 32, 111, 102, 32, 98, 111, 117, 110, 100, 115, 58, 32, 48, 32, 60, 61, 32, 116, 109, 112, 55, 32, 60, 32, 53, 48, 50, 53, 55};
+.extern .shared .align 1 .b8 global_smem[];
+
+.visible .entry triton__0d1d2de(
+	.param .u64 triton__0d1d2de_param_0,
+	.param .u64 triton__0d1d2de_param_1,
+	.param .u64 triton__0d1d2de_param_2
+)
+.maxntid 128, 1, 1
+{
+	.reg .pred 	%p<24>;
+	.reg .b16 	%rs<21>;
+	.reg .b32 	%r<21>;
+	.reg .b64 	%rd<58>;
+	.loc	1 18 0
+$L__func_begin0:
+	.loc	1 18 0
+
+	ld.param.u64 	%rd9, [triton__0d1d2de_param_1];
+	ld.param.u64 	%rd16, [triton__0d1d2de_param_0];
+$L__tmp0:
+	.loc	1 21 36
+	mov.u32 	%r4, %tid.x;
+	and.b32  	%r1, %r4, 127;
+	shl.b32 	%r2, %r1, 1;
+	or.b32  	%r5, %r2, 1;
+	or.b32  	%r6, %r2, 256;
+	.loc	1 20 28
+	mov.u32 %r3, %ctaid.x;
+	.loc	1 20 46
+	mul.wide.s32 	%rd1, %r3, 512;
+	cvt.u64.u32 	%rd17, %r2;
+	cvt.u64.u32 	%rd18, %r6;
+	.loc	1 21 23
+	or.b64  	%rd2, %rd1, %rd17;
+	or.b64  	%rd3, %rd1, %rd18;
+	.loc	1 24 30
+	shl.b64 	%rd19, %rd2, 3;
+	add.s64 	%rd12, %rd16, %rd19;
+	add.s64 	%rd15, %rd12, 2048;
+	mov.pred 	%p20, -1;
+	.loc	1 24 35
+	mov.u64 %rd10, 0x0;
+	mov.u64 %rd11, 0x0;
+	@%p20 ld.global.v2.b64 { %rd10, %rd11 }, [ %rd12 + 0 ];
+	mov.u64 %rd13, 0x0;
+	mov.u64 %rd14, 0x0;
+	@%p20 ld.global.v2.b64 { %rd13, %rd14 }, [ %rd15 + 0 ];
+	.loc	1 26 19
+	setp.eq.s64 	%p3, %rd14, -1;
+	setp.eq.s64 	%p4, %rd13, -1;
+	setp.eq.s64 	%p5, %rd11, -1;
+	setp.eq.s64 	%p6, %rd10, -1;
+	.loc	1 28 32
+	selp.b64 	%rd20, 0, %rd10, %p6;
+	selp.b64 	%rd21, 0, %rd11, %p5;
+	selp.b64 	%rd22, 0, %rd13, %p4;
+	selp.b64 	%rd23, 0, %rd14, %p3;
+	.loc	1 29 18
+	add.s64 	%rd24, %rd23, 50257;
+	add.s64 	%rd25, %rd22, 50257;
+	add.s64 	%rd26, %rd21, 50257;
+	add.s64 	%rd27, %rd20, 50257;
+	.loc	1 30 18
+	setp.lt.s64 	%p7, %rd23, 0;
+	setp.lt.s64 	%p8, %rd22, 0;
+	setp.lt.s64 	%p9, %rd21, 0;
+	setp.lt.s64 	%p10, %rd20, 0;
+	.loc	1 31 32
+	selp.b64 	%rd7, %rd27, %rd20, %p10;
+	selp.b64 	%rd6, %rd26, %rd21, %p9;
+	selp.b64 	%rd5, %rd25, %rd22, %p8;
+	selp.b64 	%rd4, %rd24, %rd23, %p7;
+	.loc	1 32 36
+	setp.lt.u64 	%p11, %rd4, 50257;
+	setp.lt.u64 	%p12, %rd5, 50257;
+	setp.lt.u64 	%p13, %rd6, 50257;
+	setp.lt.u64 	%p14, %rd7, 50257;
+	mov.u32 	%r7, global_smem;
+	add.s32 	%r8, %r7, %r2;
+	selp.u16 	%rs1, 1, 0, %p14;
+	st.shared.u8 	[%r8], %rs1;
+	cvt.u64.u32 	%rd8, %r5;
+	selp.u16 	%rs2, 1, 0, %p13;
+	st.shared.u8 	[%r8+1], %rs2;
+	bar.sync 	0;
+	add.s32 	%r9, %r7, %r1;
+	ld.shared.u8 	%rs3, [%r9];
+	ld.shared.u8 	%rs4, [%r9+128];
+	bar.sync 	0;
+	selp.u16 	%rs5, 1, 0, %p12;
+	st.shared.u8 	[%r8], %rs5;
+	selp.u16 	%rs6, 1, 0, %p11;
+	st.shared.u8 	[%r8+1], %rs6;
+	bar.sync 	0;
+	ld.shared.u8 	%rs7, [%r9];
+	ld.shared.u8 	%rs8, [%r9+128];
+	setp.eq.s16 	%p15, %rs7, 0;
+	selp.u16 	%rs9, 1, 0, %p15;
+	shl.b16 	%rs10, %rs9, 2;
+	setp.eq.s16 	%p16, %rs8, 0;
+	selp.u16 	%rs11, -1, 0, %p16;
+	shl.b16 	%rs12, %rs11, 3;
+	or.b16  	%rs13, %rs12, %rs10;
+	setp.eq.s16 	%p17, %rs4, 0;
+	selp.u16 	%rs14, 1, 0, %p17;
+	setp.eq.s16 	%p18, %rs3, 0;
+	selp.u16 	%rs15, -1, 0, %p18;
+	shl.b16 	%rs16, %rs15, 1;
+	or.b16  	%rs17, %rs14, %rs16;
+	and.b16  	%rs18, %rs17, 3;
+	or.b16  	%rs19, %rs18, %rs13;
+	.loc	1 32 51
+	and.b16  	%rs20, %rs19, 15;
+	setp.eq.s16 	%p19, %rs20, 0;
+	@%p19 bra 	$L__BB0_2;
+	mov.u64 	%rd28, assertMessage_0;
+	cvta.global.u64 	%rd29, %rd28;
+	mov.u64 	%rd30, assertFile_0;
+	cvta.global.u64 	%rd31, %rd30;
+	mov.u64 	%rd32, assertFunc_0;
+	cvta.global.u64 	%rd33, %rd32;
+	mov.b32 	%r10, 883;
+	mov.u64 	%rd34, 1;
+	{ // callseq 0, 0
+	.reg .b32 temp_param_reg;
+	.param .b64 param0;
+	st.param.b64 	[param0+0], %rd29;
+	.param .b64 param1;
+	st.param.b64 	[param1+0], %rd31;
+	.param .b32 param2;
+	st.param.b32 	[param2+0], %r10;
+	.param .b64 param3;
+	st.param.b64 	[param3+0], %rd33;
+	.param .b64 param4;
+	st.param.b64 	[param4+0], %rd34;
+	call.uni 
+	__assertfail, 
+	(
+	param0, 
+	param1, 
+	param2, 
+	param3, 
+	param4
+	);
+	} // callseq 0
+$L__BB0_2:
+	.loc	1 21 36
+	or.b32  	%r15, %r2, 257;
+	cvt.u64.u32 	%rd39, %r15;
+	.loc	1 21 23
+	or.b64  	%rd40, %rd1, %rd39;
+	or.b64  	%rd41, %rd1, %rd8;
+	.loc	1 34 25
+	shl.b64 	%rd42, %rd7, 2;
+	add.s64 	%rd43, %rd9, %rd42;
+	mul.lo.s64 	%rd44, %rd2, 201028;
+	add.s64 	%rd45, %rd43, %rd44;
+	shl.b64 	%rd46, %rd6, 2;
+	add.s64 	%rd47, %rd9, %rd46;
+	mul.lo.s64 	%rd48, %rd41, 201028;
+	add.s64 	%rd49, %rd47, %rd48;
+	shl.b64 	%rd50, %rd5, 2;
+	add.s64 	%rd51, %rd9, %rd50;
+	mul.lo.s64 	%rd52, %rd3, 201028;
+	add.s64 	%rd53, %rd51, %rd52;
+	shl.b64 	%rd54, %rd4, 2;
+	add.s64 	%rd55, %rd9, %rd54;
+	mul.lo.s64 	%rd56, %rd40, 201028;
+	add.s64 	%rd57, %rd55, %rd56;
+	.loc	1 34 51
+	bar.sync 	0;
+	shl.b32 	%r16, %r2, 3;
+	add.s32 	%r18, %r7, %r16;
+	st.shared.u64 	[%r18], %rd45;
+	st.shared.u64 	[%r18+8], %rd49;
+	bar.sync 	0;
+	shl.b32 	%r19, %r1, 3;
+	add.s32 	%r20, %r7, %r19;
+	ld.shared.u64 	%rd35, [%r20];
+	ld.shared.u64 	%rd36, [%r20+1024];
+	bar.sync 	0;
+	st.shared.u64 	[%r18], %rd53;
+	st.shared.u64 	[%r18+8], %rd57;
+	bar.sync 	0;
+	ld.shared.u64 	%rd37, [%r20];
+	ld.shared.u64 	%rd38, [%r20+1024];
+	mov.b32 	%r11, -1082130432;
+	@%p20 st.global.b32 [ %rd35 + 0 ], { %r11 };
+	@%p20 st.global.b32 [ %rd36 + 0 ], { %r11 };
+	@%p20 st.global.b32 [ %rd37 + 0 ], { %r11 };
+	@%p20 st.global.b32 [ %rd38 + 0 ], { %r11 };
+	.loc	1 34 4
+	ret;
+$L__tmp1:
+$L__func_end0:
+
+}
+	.file	1 "/tmp/torchinductor_root/hl/chlrkgpvvbdizdz7sllquet2j7zhtes6meh6kenrqxov26mswvw7.py"
+	.section	.debug_abbrev
+	{
+.b8 1
+.b8 17
+.b8 1
+.b8 37
+.b8 8
+.b8 19
+.b8 5
+.b8 3
+.b8 8
+.b8 16
+.b8 6
+.b8 27
+.b8 8
+.b8 180
+.b8 66
+.b8 12
+.b8 17
+.b8 1
+.b8 18
+.b8 1
+.b8 0
+.b8 0
+.b8 2
+.b8 46
+.b8 0
+.b8 17
+.b8 1
+.b8 18
+.b8 1
+.b8 64
+.b8 10
+.b8 135
+.b8 64
+.b8 8
+.b8 3
+.b8 8
+.b8 58
+.b8 11
+.b8 59
+.b8 11
+.b8 63
+.b8 12
+.b8 0
+.b8 0
+.b8 0
+	}
+	.section	.debug_info
+	{
+.b32 176
+.b8 2
+.b8 0
+.b32 .debug_abbrev
+.b8 8
+.b8 1
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 0
+.b8 2
+.b8 0
+.b8 99
+.b8 104
+.b8 108
+.b8 114
+.b8 107
+.b8 103
+.b8 112
+.b8 118
+.b8 118
+.b8 98
+.b8 100
+.b8 105
+.b8 122
+.b8 100
+.b8 122
+.b8 55
+.b8 115
+.b8 108
+.b8 108
+.b8 113
+.b8 117
+.b8 101
+.b8 116
+.b8 50
+.b8 106
+.b8 55
+.b8 122
+.b8 104
+.b8 116
+.b8 101
+.b8 115
+.b8 54
+.b8 109
+.b8 101
+.b8 104
+.b8 54
+.b8 107
+.b8 101
+.b8 110
+.b8 114
+.b8 113
+.b8 120
+.b8 111
+.b8 118
+.b8 50
+.b8 54
+.b8 109
+.b8 115
+.b8 119
+.b8 118
+.b8 119
+.b8 55
+.b8 46
+.b8 112
+.b8 121
+.b8 0
+.b32 .debug_line
+.b8 47
+.b8 116
+.b8 109
+.b8 112
+.b8 47
+.b8 116
+.b8 111
+.b8 114
+.b8 99
+.b8 104
+.b8 105
+.b8 110
+.b8 100
+.b8 117
+.b8 99
+.b8 116
+.b8 111
+.b8 114
+.b8 95
+.b8 114
+.b8 111
+.b8 111
+.b8 116
+.b8 47
+.b8 104
+.b8 108
+.b8 0
+.b8 1
+.b64 $L__func_begin0
+.b64 $L__func_end0
+.b8 2
+.b64 $L__func_begin0
+.b64 $L__func_end0
+.b8 1
+.b8 156
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 95
+.b8 95
+.b8 48
+.b8 100
+.b8 49
+.b8 100
+.b8 50
+.b8 100
+.b8 101
+.b8 0
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 95
+.b8 95
+.b8 48
+.b8 100
+.b8 49
+.b8 100
+.b8 50
+.b8 100
+.b8 101
+.b8 0
+.b8 1
+.b8 18
+.b8 1
+.b8 0
+	}
+	.section	.debug_pubnames
+	{
+.b32 $L__pubNames_end0-$L__pubNames_start0
+$L__pubNames_start0:
+.b8 2
+.b8 0
+.b32 .debug_info
+.b32 180
+.b32 125
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 95
+.b8 95
+.b8 48
+.b8 100
+.b8 49
+.b8 100
+.b8 50
+.b8 100
+.b8 101
+.b8 0
+.b32 0
+$L__pubNames_end0:
+	}
+	.section	.debug_pubtypes
+	{
+.b32 $L__pubTypes_end0-$L__pubTypes_start0
+$L__pubTypes_start0:
+.b8 2
+.b8 0
+.b32 .debug_info
+.b32 180
+.b32 0
+$L__pubTypes_end0:
+	}
+	.section	.debug_loc	{	}

.triton/dump/199215289adb100508718a5a762ba4d7/triton_.ttgir

@@ -0,0 +1,38 @@
+#blocked = #triton_gpu.blocked<{sizePerThread = [2], threadsPerWarp = [32], warpsPerCTA = [4], order = [0], CTAsPerCGA = [1], CTASplitNum = [1], CTAOrder = [0]}>
+#blocked1 = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0], CTAsPerCGA = [1], CTASplitNum = [1], CTAOrder = [0]}>
+module attributes {"triton_gpu.compute-capability" = 89 : i32, "triton_gpu.num-ctas" = 1 : i32, "triton_gpu.num-warps" = 4 : i32, "triton_gpu.threads-per-warp" = 32 : i32} {
+  tt.func public @triton__0d1d2de(%arg0: !tt.ptr<i64, 1> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg2: i64 {tt.divisibility = 16 : i32, tt.max_divisibility = 16 : i32}) attributes {noinline = false} {
+    %cst = arith.constant dense<50257> : tensor<512xi64, #blocked>
+    %cst_0 = arith.constant dense<0> : tensor<512xi64, #blocked>
+    %cst_1 = arith.constant dense<-1> : tensor<512xi64, #blocked>
+    %cst_2 = arith.constant dense<-1.000000e+00> : tensor<512xf32, #blocked1>
+    %c512_i64 = arith.constant 512 : i64
+    %0 = tt.get_program_id x : i32
+    %1 = arith.extsi %0 : i32 to i64
+    %2 = arith.muli %1, %c512_i64 : i64
+    %3 = tt.make_range {end = 512 : i32, start = 0 : i32} : tensor<512xi32, #blocked>
+    %4 = arith.extsi %3 : tensor<512xi32, #blocked> to tensor<512xi64, #blocked>
+    %5 = tt.splat %2 : (i64) -> tensor<512xi64, #blocked>
+    %6 = arith.addi %5, %4 : tensor<512xi64, #blocked>
+    %7 = tt.splat %arg0 : (!tt.ptr<i64, 1>) -> tensor<512x!tt.ptr<i64, 1>, #blocked>
+    %8 = tt.addptr %7, %6 : tensor<512x!tt.ptr<i64, 1>, #blocked>, tensor<512xi64, #blocked>
+    %9 = tt.load %8 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<512xi64, #blocked>
+    %10 = arith.cmpi ne, %9, %cst_1 : tensor<512xi64, #blocked>
+    %11 = arith.select %10, %9, %cst_0 : tensor<512xi1, #blocked>, tensor<512xi64, #blocked>
+    %12 = arith.addi %11, %cst : tensor<512xi64, #blocked>
+    %13 = arith.cmpi slt, %11, %cst_0 : tensor<512xi64, #blocked>
+    %14 = arith.select %13, %12, %11 : tensor<512xi1, #blocked>, tensor<512xi64, #blocked>
+    %15 = arith.cmpi sge, %14, %cst_0 : tensor<512xi64, #blocked>
+    %16 = arith.cmpi slt, %14, %cst : tensor<512xi64, #blocked>
+    %17 = arith.andi %15, %16 : tensor<512xi1, #blocked>
+    %18 = triton_gpu.convert_layout %17 : (tensor<512xi1, #blocked>) -> tensor<512xi1, #blocked1>
+    tt.assert %18, "index out of bounds: 0 <= tmp7 < 50257", "<frozen importlib._bootstrap_external>", "_call_with_frames_removed", 883 : tensor<512xi1, #blocked1>
+    %19 = arith.muli %6, %cst : tensor<512xi64, #blocked>
+    %20 = arith.addi %14, %19 : tensor<512xi64, #blocked>
+    %21 = tt.splat %arg1 : (!tt.ptr<f32, 1>) -> tensor<512x!tt.ptr<f32, 1>, #blocked>
+    %22 = tt.addptr %21, %20 : tensor<512x!tt.ptr<f32, 1>, #blocked>, tensor<512xi64, #blocked>
+    %23 = triton_gpu.convert_layout %22 : (tensor<512x!tt.ptr<f32, 1>, #blocked>) -> tensor<512x!tt.ptr<f32, 1>, #blocked1>
+    tt.store %23, %cst_2 {cache = 1 : i32, evict = 1 : i32} : tensor<512xf32, #blocked1>
+    tt.return
+  }
+}

.triton/dump/199215289adb100508718a5a762ba4d7/triton_.ttir

@@ -0,0 +1,34 @@
+module {
+  tt.func public @triton__0d1d2de(%arg0: !tt.ptr<i64, 1> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg2: i64 {tt.divisibility = 16 : i32, tt.max_divisibility = 16 : i32}) attributes {noinline = false} {
+    %cst = arith.constant dense<50257> : tensor<512xi64>
+    %cst_0 = arith.constant dense<0> : tensor<512xi64>
+    %c512_i64 = arith.constant 512 : i64
+    %cst_1 = arith.constant dense<-1.000000e+00> : tensor<512xf32>
+    %cst_2 = arith.constant dense<-1> : tensor<512xi64>
+    %0 = tt.get_program_id x : i32
+    %1 = arith.extsi %0 : i32 to i64
+    %2 = arith.muli %1, %c512_i64 : i64
+    %3 = tt.make_range {end = 512 : i32, start = 0 : i32} : tensor<512xi32>
+    %4 = arith.extsi %3 : tensor<512xi32> to tensor<512xi64>
+    %5 = tt.splat %2 : (i64) -> tensor<512xi64>
+    %6 = arith.addi %5, %4 : tensor<512xi64>
+    %7 = tt.splat %arg0 : (!tt.ptr<i64, 1>) -> tensor<512x!tt.ptr<i64, 1>>
+    %8 = tt.addptr %7, %6 : tensor<512x!tt.ptr<i64, 1>>, tensor<512xi64>
+    %9 = tt.load %8 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<512xi64>
+    %10 = arith.cmpi ne, %9, %cst_2 : tensor<512xi64>
+    %11 = arith.select %10, %9, %cst_0 : tensor<512xi1>, tensor<512xi64>
+    %12 = arith.addi %11, %cst : tensor<512xi64>
+    %13 = arith.cmpi slt, %11, %cst_0 : tensor<512xi64>
+    %14 = arith.select %13, %12, %11 : tensor<512xi1>, tensor<512xi64>
+    %15 = arith.cmpi sge, %14, %cst_0 : tensor<512xi64>
+    %16 = arith.cmpi slt, %14, %cst : tensor<512xi64>
+    %17 = arith.andi %15, %16 : tensor<512xi1>
+    tt.assert %17, "index out of bounds: 0 <= tmp7 < 50257", "<frozen importlib._bootstrap_external>", "_call_with_frames_removed", 883 : tensor<512xi1>
+    %18 = arith.muli %6, %cst : tensor<512xi64>
+    %19 = arith.addi %14, %18 : tensor<512xi64>
+    %20 = tt.splat %arg1 : (!tt.ptr<f32, 1>) -> tensor<512x!tt.ptr<f32, 1>>
+    %21 = tt.addptr %20, %19 : tensor<512x!tt.ptr<f32, 1>>, tensor<512xi64>
+    tt.store %21, %cst_1 {cache = 1 : i32, evict = 1 : i32} : tensor<512xf32>
+    tt.return
+  }
+}

.triton/dump/1c14bdb6903aa6825e214bbdf57fd077/triton_.ptx

@@ -0,0 +1,312 @@
+//
+// Generated by LLVM NVPTX Back-End
+//
+
+.version 8.2
+.target sm_89
+.address_size 64
+
+	// .globl	triton__0d1d2de
+
+.visible .entry triton__0d1d2de(
+	.param .u64 triton__0d1d2de_param_0,
+	.param .u64 triton__0d1d2de_param_1,
+	.param .u32 triton__0d1d2de_param_2
+)
+.maxntid 128, 1, 1
+{
+	.reg .pred 	%p<4>;
+	.reg .b16 	%rs<9>;
+	.reg .b32 	%r<31>;
+	.reg .b64 	%rd<8>;
+	.loc	1 18 0
+$L__func_begin0:
+	.loc	1 18 0
+
+	ld.param.u64 	%rd4, [triton__0d1d2de_param_0];
+	ld.param.u64 	%rd5, [triton__0d1d2de_param_1];
+$L__tmp0:
+	.loc	1 21 36
+	mov.u32 	%r22, %tid.x;
+	shl.b32 	%r23, %r22, 3;
+	and.b32  	%r24, %r23, 1016;
+	.loc	1 20 28
+	mov.u32 %r1, %ctaid.x;
+	.loc	1 20 33
+	shl.b32 	%r25, %r1, 10;
+	.loc	1 21 23
+	or.b32  	%r26, %r25, %r24;
+	.loc	1 24 30
+	mul.wide.s32 	%rd6, %r26, 4;
+	add.s64 	%rd1, %rd4, %rd6;
+	add.s64 	%rd2, %rd1, 16;
+	mov.pred 	%p1, -1;
+	.loc	1 24 35
+	mov.u32 %r10, 0x0;
+	mov.u32 %r11, 0x0;
+	mov.u32 %r12, 0x0;
+	mov.u32 %r13, 0x0;
+	@%p1 ld.global.v4.b32 { %r10, %r11, %r12, %r13 }, [ %rd1 + 0 ];
+	mov.u32 %r14, 0x0;
+	mov.u32 %r15, 0x0;
+	mov.u32 %r16, 0x0;
+	mov.u32 %r17, 0x0;
+	@%p1 ld.global.v4.b32 { %r14, %r15, %r16, %r17 }, [ %rd2 + 0 ];
+	.loc	1 26 25
+	mul.wide.s32 	%rd7, %r26, 2;
+	add.s64 	%rd3, %rd5, %rd7;
+	.loc	1 26 36
+	cvt.rn.bf16.f32 %rs1, %r10;
+	cvt.rn.bf16.f32 %rs2, %r11;
+	cvt.rn.bf16.f32 %rs3, %r12;
+	cvt.rn.bf16.f32 %rs4, %r13;
+	cvt.rn.bf16.f32 %rs5, %r14;
+	cvt.rn.bf16.f32 %rs6, %r15;
+	cvt.rn.bf16.f32 %rs7, %r16;
+	cvt.rn.bf16.f32 %rs8, %r17;
+	mov.b32 	%r27, {%rs1, %rs2};
+	mov.b32 	%r28, {%rs3, %rs4};
+	mov.b32 	%r29, {%rs5, %rs6};
+	mov.b32 	%r30, {%rs7, %rs8};
+	@%p1 st.global.v4.b32 [ %rd3 + 0 ], { %r27, %r28, %r29, %r30 };
+	.loc	1 26 4
+	ret;
+$L__tmp1:
+$L__func_end0:
+
+}
+	.file	1 "/tmp/torchinductor_root/5t/c5tryp5qwkhreijk7s5x327wofz54lwj4kvctuqdzv2vrf2xyons.py"
+	.section	.debug_abbrev
+	{
+.b8 1
+.b8 17
+.b8 1
+.b8 37
+.b8 8
+.b8 19
+.b8 5
+.b8 3
+.b8 8
+.b8 16
+.b8 6
+.b8 27
+.b8 8
+.b8 180
+.b8 66
+.b8 12
+.b8 17
+.b8 1
+.b8 18
+.b8 1
+.b8 0
+.b8 0
+.b8 2
+.b8 46
+.b8 0
+.b8 17
+.b8 1
+.b8 18
+.b8 1
+.b8 64
+.b8 10
+.b8 135
+.b8 64
+.b8 8
+.b8 3
+.b8 8
+.b8 58
+.b8 11
+.b8 59
+.b8 11
+.b8 63
+.b8 12
+.b8 0
+.b8 0
+.b8 0
+	}
+	.section	.debug_info
+	{
+.b32 176
+.b8 2
+.b8 0
+.b32 .debug_abbrev
+.b8 8
+.b8 1
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 0
+.b8 2
+.b8 0
+.b8 99
+.b8 53
+.b8 116
+.b8 114
+.b8 121
+.b8 112
+.b8 53
+.b8 113
+.b8 119
+.b8 107
+.b8 104
+.b8 114
+.b8 101
+.b8 105
+.b8 106
+.b8 107
+.b8 55
+.b8 115
+.b8 53
+.b8 120
+.b8 51
+.b8 50
+.b8 55
+.b8 119
+.b8 111
+.b8 102
+.b8 122
+.b8 53
+.b8 52
+.b8 108
+.b8 119
+.b8 106
+.b8 52
+.b8 107
+.b8 118
+.b8 99
+.b8 116
+.b8 117
+.b8 113
+.b8 100
+.b8 122
+.b8 118
+.b8 50
+.b8 118
+.b8 114
+.b8 102
+.b8 50
+.b8 120
+.b8 121
+.b8 111
+.b8 110
+.b8 115
+.b8 46
+.b8 112
+.b8 121
+.b8 0
+.b32 .debug_line
+.b8 47
+.b8 116
+.b8 109
+.b8 112
+.b8 47
+.b8 116
+.b8 111
+.b8 114
+.b8 99
+.b8 104
+.b8 105
+.b8 110
+.b8 100
+.b8 117
+.b8 99
+.b8 116
+.b8 111
+.b8 114
+.b8 95
+.b8 114
+.b8 111
+.b8 111
+.b8 116
+.b8 47
+.b8 53
+.b8 116
+.b8 0
+.b8 1
+.b64 $L__func_begin0
+.b64 $L__func_end0
+.b8 2
+.b64 $L__func_begin0
+.b64 $L__func_end0
+.b8 1
+.b8 156
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 95
+.b8 95
+.b8 48
+.b8 100
+.b8 49
+.b8 100
+.b8 50
+.b8 100
+.b8 101
+.b8 0
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 95
+.b8 95
+.b8 48
+.b8 100
+.b8 49
+.b8 100
+.b8 50
+.b8 100
+.b8 101
+.b8 0
+.b8 1
+.b8 18
+.b8 1
+.b8 0
+	}
+	.section	.debug_pubnames
+	{
+.b32 $L__pubNames_end0-$L__pubNames_start0
+$L__pubNames_start0:
+.b8 2
+.b8 0
+.b32 .debug_info
+.b32 180
+.b32 125
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 95
+.b8 95
+.b8 48
+.b8 100
+.b8 49
+.b8 100
+.b8 50
+.b8 100
+.b8 101
+.b8 0
+.b32 0
+$L__pubNames_end0:
+	}
+	.section	.debug_pubtypes
+	{
+.b32 $L__pubTypes_end0-$L__pubTypes_start0
+$L__pubTypes_start0:
+.b8 2
+.b8 0
+.b32 .debug_info
+.b32 180
+.b32 0
+$L__pubTypes_end0:
+	}
+	.section	.debug_loc	{	}

.triton/dump/1c14bdb6903aa6825e214bbdf57fd077/triton_.ttgir

@@ -0,0 +1,19 @@
+#blocked = #triton_gpu.blocked<{sizePerThread = [8], threadsPerWarp = [32], warpsPerCTA = [4], order = [0], CTAsPerCGA = [1], CTASplitNum = [1], CTAOrder = [0]}>
+module attributes {"triton_gpu.compute-capability" = 89 : i32, "triton_gpu.num-ctas" = 1 : i32, "triton_gpu.num-warps" = 4 : i32, "triton_gpu.threads-per-warp" = 32 : i32} {
+  tt.func public @triton__0d1d2de(%arg0: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<bf16, 1> {tt.divisibility = 16 : i32}, %arg2: i32 {tt.divisibility = 16 : i32, tt.max_divisibility = 16 : i32}) attributes {noinline = false} {
+    %c1024_i32 = arith.constant 1024 : i32
+    %0 = tt.get_program_id x : i32
+    %1 = arith.muli %0, %c1024_i32 : i32
+    %2 = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32, #blocked>
+    %3 = tt.splat %1 : (i32) -> tensor<1024xi32, #blocked>
+    %4 = arith.addi %3, %2 : tensor<1024xi32, #blocked>
+    %5 = tt.splat %arg0 : (!tt.ptr<f32, 1>) -> tensor<1024x!tt.ptr<f32, 1>, #blocked>
+    %6 = tt.addptr %5, %4 : tensor<1024x!tt.ptr<f32, 1>, #blocked>, tensor<1024xi32, #blocked>
+    %7 = tt.load %6 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<1024xf32, #blocked>
+    %8 = tt.splat %arg1 : (!tt.ptr<bf16, 1>) -> tensor<1024x!tt.ptr<bf16, 1>, #blocked>
+    %9 = tt.addptr %8, %4 : tensor<1024x!tt.ptr<bf16, 1>, #blocked>, tensor<1024xi32, #blocked>
+    %10 = arith.truncf %7 : tensor<1024xf32, #blocked> to tensor<1024xbf16, #blocked>
+    tt.store %9, %10 {cache = 1 : i32, evict = 1 : i32} : tensor<1024xbf16, #blocked>
+    tt.return
+  }
+}

.triton/dump/1c188b233fcb854770e6a3cf1802c844/triton_.ttir

@@ -0,0 +1,56 @@
+module {
+  tt.func public @triton__0d1d2d3de4de(%arg0: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg3: i32 {tt.divisibility = 16 : i32, tt.max_divisibility = 16 : i32}, %arg4: i32 {tt.divisibility = 16 : i32, tt.max_divisibility = 16 : i32}) attributes {noinline = false} {
+    %c8_i32 = arith.constant 8 : i32
+    %c128_i32 = arith.constant 128 : i32
+    %c0_i32 = arith.constant 0 : i32
+    %cst = arith.constant dense<32768> : tensor<64x1xi32>
+    %cst_0 = arith.constant dense<256> : tensor<1x8xi32>
+    %cst_1 = arith.constant dense<128> : tensor<1x8xi32>
+    %cst_2 = arith.constant dense<0.000000e+00> : tensor<64x8xf32>
+    %cst_3 = arith.constant dense<256> : tensor<64x1xi32>
+    %c64_i32 = arith.constant 64 : i32
+    %0 = tt.get_program_id x : i32
+    %1 = arith.muli %0, %c64_i32 : i32
+    %2 = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32>
+    %3 = tt.expand_dims %2 {axis = 1 : i32} : (tensor<64xi32>) -> tensor<64x1xi32>
+    %4 = tt.splat %1 : (i32) -> tensor<64x1xi32>
+    %5 = arith.addi %4, %3 : tensor<64x1xi32>
+    %6 = tt.make_range {end = 8 : i32, start = 0 : i32} : tensor<8xi32>
+    %7 = tt.expand_dims %6 {axis = 0 : i32} : (tensor<8xi32>) -> tensor<1x8xi32>
+    %8 = arith.remsi %5, %cst_3 : tensor<64x1xi32>
+    %9 = arith.divsi %5, %cst_3 : tensor<64x1xi32>
+    %10 = tt.broadcast %8 : (tensor<64x1xi32>) -> tensor<64x8xi32>
+    %11 = arith.muli %9, %cst : tensor<64x1xi32>
+    %12 = tt.broadcast %11 : (tensor<64x1xi32>) -> tensor<64x8xi32>
+    %13 = tt.splat %arg0 : (!tt.ptr<f32, 1>) -> tensor<64x8x!tt.ptr<f32, 1>>
+    %14 = tt.splat %arg1 : (!tt.ptr<f32, 1>) -> tensor<64x8x!tt.ptr<f32, 1>>
+    %15 = scf.for %arg5 = %c0_i32 to %c128_i32 step %c8_i32 iter_args(%arg6 = %cst_2) -> (tensor<64x8xf32>)  : i32 {
+      %20 = tt.splat %arg5 : (i32) -> tensor<1x8xi32>
+      %21 = arith.addi %20, %7 : tensor<1x8xi32>
+      %22 = arith.cmpi slt, %21, %cst_1 : tensor<1x8xi32>
+      %23 = arith.muli %21, %cst_0 : tensor<1x8xi32>
+      %24 = tt.broadcast %23 : (tensor<1x8xi32>) -> tensor<64x8xi32>
+      %25 = arith.addi %10, %24 : tensor<64x8xi32>
+      %26 = arith.addi %25, %12 : tensor<64x8xi32>
+      %27 = tt.addptr %13, %26 : tensor<64x8x!tt.ptr<f32, 1>>, tensor<64x8xi32>
+      %28 = tt.broadcast %22 : (tensor<1x8xi1>) -> tensor<64x8xi1>
+      %29 = tt.load %27, %28, %cst_2 {cache = 1 : i32, evict = 2 : i32, isVolatile = false} : tensor<64x8xf32>
+      %30 = tt.addptr %14, %26 : tensor<64x8x!tt.ptr<f32, 1>>, tensor<64x8xi32>
+      %31 = tt.load %30, %28, %cst_2 {cache = 1 : i32, evict = 2 : i32, isVolatile = false} : tensor<64x8xf32>
+      %32 = arith.mulf %29, %31 : tensor<64x8xf32>
+      %33 = arith.addf %arg6, %32 : tensor<64x8xf32>
+      %34 = arith.select %28, %33, %arg6 : tensor<64x8xi1>, tensor<64x8xf32>
+      scf.yield %34 : tensor<64x8xf32>
+    }
+    %16 = "tt.reduce"(%15) <{axis = 1 : i32}> ({
+    ^bb0(%arg5: f32, %arg6: f32):
+      %20 = arith.addf %arg5, %arg6 : f32
+      tt.reduce.return %20 : f32
+    }) : (tensor<64x8xf32>) -> tensor<64xf32>
+    %17 = tt.expand_dims %16 {axis = 1 : i32} : (tensor<64xf32>) -> tensor<64x1xf32>
+    %18 = tt.splat %arg2 : (!tt.ptr<f32, 1>) -> tensor<64x1x!tt.ptr<f32, 1>>
+    %19 = tt.addptr %18, %5 : tensor<64x1x!tt.ptr<f32, 1>>, tensor<64x1xi32>
+    tt.store %19, %17 {cache = 1 : i32, evict = 1 : i32} : tensor<64x1xf32>
+    tt.return
+  }
+}

.triton/dump/21d0195c63fb062bfc567b79c9bb2771/triton_.ttgir

@@ -0,0 +1,88 @@
+#blocked = #triton_gpu.blocked<{sizePerThread = [4], threadsPerWarp = [32], warpsPerCTA = [2], order = [0], CTAsPerCGA = [1], CTASplitNum = [1], CTAOrder = [0]}>
+module attributes {"triton_gpu.compute-capability" = 89 : i32, "triton_gpu.num-ctas" = 1 : i32, "triton_gpu.num-warps" = 2 : i32, "triton_gpu.threads-per-warp" = 32 : i32} {
+  tt.func public @triton__0d1d2d3d4d5d6d7d8de9de(%arg0: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<bf16, 1> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg3: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg4: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg5: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg6: !tt.ptr<i64, 1> {tt.divisibility = 16 : i32}, %arg7: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg8: i32 {tt.divisibility = 16 : i32, tt.max_divisibility = 16 : i32}, %arg9: i32 {tt.divisibility = 16 : i32, tt.max_divisibility = 16 : i32}) attributes {noinline = false} {
+    %cst = arith.constant dense<256> : tensor<256xi32, #blocked>
+    %cst_0 = arith.constant dense<-1> : tensor<1xi64, #blocked>
+    %cst_1 = arith.constant dense<2.560000e+02> : tensor<1xf32, #blocked>
+    %cst_2 = arith.constant dense<256> : tensor<1xi64, #blocked>
+    %cst_3 = arith.constant dense<0> : tensor<1xi64, #blocked>
+    %cst_4 = arith.constant dense<50257> : tensor<1xi64, #blocked>
+    %cst_5 = arith.constant 0.000000e+00 : f32
+    %c256_i32 = arith.constant 256 : i32
+    %cst_6 = arith.constant dense<0.000000e+00> : tensor<256xf32, #blocked>
+    %cst_7 = arith.constant dense<2.560000e+02> : tensor<256xf32, #blocked>
+    %cst_8 = arith.constant dense<0.000000e+00> : tensor<256xbf16, #blocked>
+    %0 = tt.get_program_id x : i32
+    %1 = tt.make_range {end = 256 : i32, start = 0 : i32} : tensor<256xi32, #blocked>
+    %2 = arith.cmpi slt, %1, %cst : tensor<256xi32, #blocked>
+    %3 = arith.muli %0, %c256_i32 : i32
+    %4 = tt.splat %3 : (i32) -> tensor<256xi32, #blocked>
+    %5 = arith.addi %1, %4 : tensor<256xi32, #blocked>
+    %6 = tt.splat %arg1 : (!tt.ptr<bf16, 1>) -> tensor<256x!tt.ptr<bf16, 1>, #blocked>
+    %7 = tt.addptr %6, %5 : tensor<256x!tt.ptr<bf16, 1>, #blocked>, tensor<256xi32, #blocked>
+    %8 = tt.load %7, %2, %cst_8 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<256xbf16, #blocked>
+    %9 = arith.extf %8 : tensor<256xbf16, #blocked> to tensor<256xf32, #blocked>
+    %10 = tt.splat %arg2 : (!tt.ptr<f32, 1>) -> tensor<256x!tt.ptr<f32, 1>, #blocked>
+    %11 = tt.addptr %10, %1 : tensor<256x!tt.ptr<f32, 1>, #blocked>, tensor<256xi32, #blocked>
+    %12 = tt.load %11, %2, %cst_6 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<256xf32, #blocked>
+    %13 = tt.splat %arg3 : (!tt.ptr<f32, 1>) -> tensor<256x!tt.ptr<f32, 1>, #blocked>
+    %14 = tt.addptr %13, %5 : tensor<256x!tt.ptr<f32, 1>, #blocked>, tensor<256xi32, #blocked>
+    %15 = tt.load %14, %2, %cst_6 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<256xf32, #blocked>
+    %16 = tt.addptr %arg4, %0 : !tt.ptr<f32, 1>, i32
+    %17 = tt.splat %16 : (!tt.ptr<f32, 1>) -> tensor<1x!tt.ptr<f32, 1>, #blocked>
+    %18 = tt.load %17 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<1xf32, #blocked>
+    %19 = tt.addptr %arg5, %0 : !tt.ptr<f32, 1>, i32
+    %20 = tt.splat %19 : (!tt.ptr<f32, 1>) -> tensor<1x!tt.ptr<f32, 1>, #blocked>
+    %21 = tt.load %20 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<1xf32, #blocked>
+    %22 = tt.addptr %arg6, %0 : !tt.ptr<i64, 1>, i32
+    %23 = tt.splat %22 : (!tt.ptr<i64, 1>) -> tensor<1x!tt.ptr<i64, 1>, #blocked>
+    %24 = tt.load %23 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<1xi64, #blocked>
+    %25 = tt.splat %arg0 : (!tt.ptr<f32, 1>) -> tensor<256x!tt.ptr<f32, 1>, #blocked>
+    %26 = tt.addptr %25, %5 : tensor<256x!tt.ptr<f32, 1>, #blocked>, tensor<256xi32, #blocked>
+    %27 = tt.load %26, %2, %cst_6 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<256xf32, #blocked>
+    %28 = arith.mulf %9, %12 : tensor<256xf32, #blocked>
+    %29 = arith.select %2, %28, %cst_6 : tensor<256xi1, #blocked>, tensor<256xf32, #blocked>
+    %30 = "tt.reduce"(%29) <{axis = 0 : i32}> ({
+    ^bb0(%arg10: f32, %arg11: f32):
+      %63 = arith.addf %arg10, %arg11 : f32
+      tt.reduce.return %63 : f32
+    }) : (tensor<256xf32, #blocked>) -> f32
+    %31 = arith.addf %30, %cst_5 : f32
+    %32 = tt.broadcast %18 : (tensor<1xf32, #blocked>) -> tensor<256xf32, #blocked>
+    %33 = arith.subf %15, %32 : tensor<256xf32, #blocked>
+    %34 = tt.broadcast %21 : (tensor<1xf32, #blocked>) -> tensor<256xf32, #blocked>
+    %35 = arith.mulf %33, %34 : tensor<256xf32, #blocked>
+    %36 = arith.mulf %28, %35 : tensor<256xf32, #blocked>
+    %37 = arith.select %2, %36, %cst_6 : tensor<256xi1, #blocked>, tensor<256xf32, #blocked>
+    %38 = "tt.reduce"(%37) <{axis = 0 : i32}> ({
+    ^bb0(%arg10: f32, %arg11: f32):
+      %63 = arith.addf %arg10, %arg11 : f32
+      tt.reduce.return %63 : f32
+    }) : (tensor<256xf32, #blocked>) -> f32
+    %39 = arith.addf %38, %cst_5 : f32
+    %40 = arith.cmpi eq, %24, %cst_0 : tensor<1xi64, #blocked>
+    %41 = arith.divf %21, %cst_1 : tensor<1xf32, #blocked>
+    %42 = arith.mulf %28, %cst_7 : tensor<256xf32, #blocked>
+    %43 = tt.splat %31 : (f32) -> tensor<256xf32, #blocked>
+    %44 = arith.subf %42, %43 : tensor<256xf32, #blocked>
+    %45 = tt.splat %39 : (f32) -> tensor<256xf32, #blocked>
+    %46 = arith.mulf %35, %45 : tensor<256xf32, #blocked>
+    %47 = arith.subf %44, %46 : tensor<256xf32, #blocked>
+    %48 = tt.broadcast %41 : (tensor<1xf32, #blocked>) -> tensor<256xf32, #blocked>
+    %49 = arith.mulf %48, %47 : tensor<256xf32, #blocked>
+    %50 = arith.addf %27, %49 : tensor<256xf32, #blocked>
+    %51 = tt.broadcast %40 : (tensor<1xi1, #blocked>) -> tensor<256xi1, #blocked>
+    %52 = arith.select %51, %cst_6, %50 : tensor<256xi1, #blocked>, tensor<256xf32, #blocked>
+    %53 = arith.addi %24, %cst_4 : tensor<1xi64, #blocked>
+    %54 = arith.cmpi slt, %24, %cst_3 : tensor<1xi64, #blocked>
+    %55 = arith.select %54, %53, %24 : tensor<1xi1, #blocked>, tensor<1xi64, #blocked>
+    %56 = arith.muli %55, %cst_2 : tensor<1xi64, #blocked>
+    %57 = tt.broadcast %56 : (tensor<1xi64, #blocked>) -> tensor<256xi64, #blocked>
+    %58 = arith.extsi %1 : tensor<256xi32, #blocked> to tensor<256xi64, #blocked>
+    %59 = arith.addi %58, %57 : tensor<256xi64, #blocked>
+    %60 = tt.splat %arg7 : (!tt.ptr<f32, 1>) -> tensor<256x!tt.ptr<f32, 1>, #blocked>
+    %61 = tt.addptr %60, %59 : tensor<256x!tt.ptr<f32, 1>, #blocked>, tensor<256xi64, #blocked>
+    %62 = "tt.atomic_rmw"(%61, %52, %2) <{atomic_rmw_op = 5 : i32, scope = 1 : i32, sem = 4 : i32}> : (tensor<256x!tt.ptr<f32, 1>, #blocked>, tensor<256xf32, #blocked>, tensor<256xi1, #blocked>) -> tensor<256xf32, #blocked>
+    tt.return
+  }
+}

.triton/dump/415aac87553b7d064f52694fa7254686/triton_.ptx

@@ -0,0 +1,778 @@
+//
+// Generated by LLVM NVPTX Back-End
+//
+
+.version 8.2
+.target sm_89
+.address_size 64
+
+	// .globl	triton__0d1d2de
+.global .align 1 .b8 _$_str[11] = {95, 95, 67, 85, 68, 65, 95, 70, 84, 90, 0};
+
+.visible .entry triton__0d1d2de(
+	.param .u64 triton__0d1d2de_param_0,
+	.param .u64 triton__0d1d2de_param_1,
+	.param .u32 triton__0d1d2de_param_2
+)
+.maxntid 128, 1, 1
+{
+	.reg .pred 	%p<27>;
+	.reg .b16 	%rs<17>;
+	.reg .b32 	%r<67>;
+	.reg .f32 	%f<431>;
+	.reg .b64 	%rd<8>;
+	.loc	1 18 0
+$L__func_begin0:
+	.loc	1 18 0
+
+	ld.param.u64 	%rd4, [triton__0d1d2de_param_0];
+$L__tmp0:
+	.loc	1 21 36
+	mov.u32 	%r14, %tid.x;
+	shl.b32 	%r15, %r14, 3;
+	and.b32  	%r16, %r15, 1016;
+	.loc	1 20 28
+	mov.u32 %r1, %ctaid.x;
+	.loc	1 20 33
+	shl.b32 	%r17, %r1, 10;
+	.loc	1 21 23
+	or.b32  	%r18, %r17, %r16;
+	.loc	1 24 30
+	mul.wide.s32 	%rd5, %r18, 2;
+	add.s64 	%rd3, %rd4, %rd5;
+	mov.pred 	%p1, -1;
+	.loc	1 24 35
+	mov.u32 %r2, 0x0;
+	mov.u32 %r3, 0x0;
+	mov.u32 %r4, 0x0;
+	mov.u32 %r5, 0x0;
+	@%p1 ld.global.v4.b32 { %r2, %r3, %r4, %r5 }, [ %rd3 + 0 ];
+	cvt.u16.u32 	%rs1, %r2;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs2}, %r2; }
+	cvt.u16.u32 	%rs3, %r3;
+	.loc	1 24 44
+	cvt.f32.bf16 %r6, %rs1;
+	mov.b32 	%f1, %r6;
+	cvt.f32.bf16 %r7, %rs2;
+	mov.b32 	%f2, %r7;
+	.loc	1 29 18
+	mul.f32 	%f9, %f1, 0f3F3504F3;
+	.loc	1 30 23
+	abs.ftz.f32 	%f17, %f9;
+	setp.ge.f32 	%p2, %f17, 0f3F8060FE;
+	mov.f32 	%f365, 0f3789CA3C;
+	mov.f32 	%f364, 0fB9F560B9;
+	mov.f32 	%f363, 0f3BAC840B;
+	mov.f32 	%f362, 0fBD0C8162;
+	mov.f32 	%f361, 0f3E1CF906;
+	mov.f32 	%f360, 0f3F6A937E;
+	mov.f32 	%f359, 0f3F20D842;
+	mov.f32 	%f366, %f17;
+	@%p2 bra 	$L__BB0_2;
+	.loc	1 0 23
+	mov.f32 	%f365, 0f38B1E96A;
+	mov.f32 	%f364, 0fBA574D20;
+	mov.f32 	%f363, 0f3BAAD5EA;
+	mov.f32 	%f362, 0fBCDC1BE7;
+	mov.f32 	%f361, 0f3DE718AF;
+	mov.f32 	%f360, 0fBEC093AC;
+	mov.f32 	%f359, 0f3E0375D3;
+	.loc	1 30 23
+	mul.f32 	%f366, %f9, %f9;
+$L__BB0_2:
+	.loc	1 0 0
+	cvt.f32.bf16 %r8, %rs3;
+	mul.f32 	%f10, %f2, 0f3F3504F3;
+	.loc	1 30 23
+	setp.ltu.f32 	%p3, %f17, 0f3F8060FE;
+	fma.rn.ftz.f32 	%f135, %f365, %f366, %f364;
+	fma.rn.ftz.f32 	%f136, %f135, %f366, %f363;
+	fma.rn.ftz.f32 	%f137, %f136, %f366, %f362;
+	fma.rn.ftz.f32 	%f138, %f137, %f366, %f361;
+	fma.rn.ftz.f32 	%f139, %f138, %f366, %f360;
+	fma.rn.ftz.f32 	%f140, %f139, %f366, %f359;
+	neg.f32 	%f141, %f366;
+	selp.f32 	%f142, %f141, %f9, %p2;
+	fma.rn.ftz.f32 	%f367, %f140, %f142, %f142;
+	mov.f32 	%f358, 0f3F800000;
+	@%p3 bra 	$L__BB0_4;
+	ex2.approx.ftz.f32 	%f143, %f367;
+	sub.f32 	%f145, %f358, %f143;
+	mov.b32 	%r19, %f145;
+	mov.b32 	%r20, %f9;
+	and.b32  	%r21, %r20, -2147483648;
+	or.b32  	%r22, %r21, %r19;
+	mov.b32 	%f367, %r22;
+$L__BB0_4:
+	.loc	1 0 0
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs4}, %r3; }
+	mov.b32 	%f3, %r8;
+	.loc	1 30 23
+	abs.ftz.f32 	%f30, %f10;
+	setp.ge.f32 	%p5, %f30, 0f3F8060FE;
+	mov.f32 	%f374, 0f3789CA3C;
+	mov.f32 	%f373, 0fB9F560B9;
+	mov.f32 	%f372, 0f3BAC840B;
+	mov.f32 	%f371, 0fBD0C8162;
+	mov.f32 	%f370, 0f3E1CF906;
+	mov.f32 	%f369, 0f3F6A937E;
+	mov.f32 	%f368, 0f3F20D842;
+	mov.f32 	%f375, %f30;
+	@%p5 bra 	$L__BB0_6;
+	mul.f32 	%f375, %f10, %f10;
+	mov.f32 	%f374, 0f38B1E96A;
+	mov.f32 	%f373, 0fBA574D20;
+	mov.f32 	%f372, 0f3BAAD5EA;
+	mov.f32 	%f371, 0fBCDC1BE7;
+	mov.f32 	%f370, 0f3DE718AF;
+	mov.f32 	%f369, 0fBEC093AC;
+	mov.f32 	%f368, 0f3E0375D3;
+$L__BB0_6:
+	.loc	1 0 0
+	cvt.f32.bf16 %r9, %rs4;
+	mul.f32 	%f11, %f3, 0f3F3504F3;
+	.loc	1 30 23
+	setp.ltu.f32 	%p6, %f30, 0f3F8060FE;
+	fma.rn.ftz.f32 	%f160, %f374, %f375, %f373;
+	fma.rn.ftz.f32 	%f161, %f160, %f375, %f372;
+	fma.rn.ftz.f32 	%f162, %f161, %f375, %f371;
+	fma.rn.ftz.f32 	%f163, %f162, %f375, %f370;
+	fma.rn.ftz.f32 	%f164, %f163, %f375, %f369;
+	fma.rn.ftz.f32 	%f165, %f164, %f375, %f368;
+	neg.f32 	%f166, %f375;
+	selp.f32 	%f167, %f166, %f10, %p5;
+	fma.rn.ftz.f32 	%f376, %f165, %f167, %f167;
+	@%p6 bra 	$L__BB0_8;
+	ex2.approx.ftz.f32 	%f168, %f376;
+	sub.f32 	%f170, %f358, %f168;
+	mov.b32 	%r23, %f170;
+	mov.b32 	%r24, %f10;
+	and.b32  	%r25, %r24, -2147483648;
+	or.b32  	%r26, %r25, %r23;
+	mov.b32 	%f376, %r26;
+$L__BB0_8:
+	.loc	1 0 0
+	cvt.u16.u32 	%rs5, %r4;
+	mov.b32 	%f4, %r9;
+	.loc	1 30 23
+	abs.ftz.f32 	%f43, %f11;
+	setp.ge.f32 	%p8, %f43, 0f3F8060FE;
+	mov.f32 	%f383, 0f3789CA3C;
+	mov.f32 	%f382, 0fB9F560B9;
+	mov.f32 	%f381, 0f3BAC840B;
+	mov.f32 	%f380, 0fBD0C8162;
+	mov.f32 	%f379, 0f3E1CF906;
+	mov.f32 	%f378, 0f3F6A937E;
+	mov.f32 	%f377, 0f3F20D842;
+	mov.f32 	%f384, %f43;
+	@%p8 bra 	$L__BB0_10;
+	mul.f32 	%f384, %f11, %f11;
+	mov.f32 	%f383, 0f38B1E96A;
+	mov.f32 	%f382, 0fBA574D20;
+	mov.f32 	%f381, 0f3BAAD5EA;
+	mov.f32 	%f380, 0fBCDC1BE7;
+	mov.f32 	%f379, 0f3DE718AF;
+	mov.f32 	%f378, 0fBEC093AC;
+	mov.f32 	%f377, 0f3E0375D3;
+$L__BB0_10:
+	.loc	1 0 0
+	cvt.f32.bf16 %r10, %rs5;
+	mul.f32 	%f12, %f4, 0f3F3504F3;
+	.loc	1 30 23
+	setp.ltu.f32 	%p9, %f43, 0f3F8060FE;
+	fma.rn.ftz.f32 	%f185, %f383, %f384, %f382;
+	fma.rn.ftz.f32 	%f186, %f185, %f384, %f381;
+	fma.rn.ftz.f32 	%f187, %f186, %f384, %f380;
+	fma.rn.ftz.f32 	%f188, %f187, %f384, %f379;
+	fma.rn.ftz.f32 	%f189, %f188, %f384, %f378;
+	fma.rn.ftz.f32 	%f190, %f189, %f384, %f377;
+	neg.f32 	%f191, %f384;
+	selp.f32 	%f192, %f191, %f11, %p8;
+	fma.rn.ftz.f32 	%f385, %f190, %f192, %f192;
+	@%p9 bra 	$L__BB0_12;
+	ex2.approx.ftz.f32 	%f193, %f385;
+	sub.f32 	%f195, %f358, %f193;
+	mov.b32 	%r27, %f195;
+	mov.b32 	%r28, %f11;
+	and.b32  	%r29, %r28, -2147483648;
+	or.b32  	%r30, %r29, %r27;
+	mov.b32 	%f385, %r30;
+$L__BB0_12:
+	.loc	1 0 0
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs6}, %r4; }
+	mov.b32 	%f5, %r10;
+	.loc	1 30 23
+	abs.ftz.f32 	%f56, %f12;
+	setp.ge.f32 	%p11, %f56, 0f3F8060FE;
+	mov.f32 	%f392, 0f3789CA3C;
+	mov.f32 	%f391, 0fB9F560B9;
+	mov.f32 	%f390, 0f3BAC840B;
+	mov.f32 	%f389, 0fBD0C8162;
+	mov.f32 	%f388, 0f3E1CF906;
+	mov.f32 	%f387, 0f3F6A937E;
+	mov.f32 	%f386, 0f3F20D842;
+	mov.f32 	%f393, %f56;
+	@%p11 bra 	$L__BB0_14;
+	mul.f32 	%f393, %f12, %f12;
+	mov.f32 	%f392, 0f38B1E96A;
+	mov.f32 	%f391, 0fBA574D20;
+	mov.f32 	%f390, 0f3BAAD5EA;
+	mov.f32 	%f389, 0fBCDC1BE7;
+	mov.f32 	%f388, 0f3DE718AF;
+	mov.f32 	%f387, 0fBEC093AC;
+	mov.f32 	%f386, 0f3E0375D3;
+$L__BB0_14:
+	.loc	1 0 0
+	cvt.f32.bf16 %r11, %rs6;
+	mul.f32 	%f13, %f5, 0f3F3504F3;
+	.loc	1 30 23
+	setp.ltu.f32 	%p12, %f56, 0f3F8060FE;
+	fma.rn.ftz.f32 	%f210, %f392, %f393, %f391;
+	fma.rn.ftz.f32 	%f211, %f210, %f393, %f390;
+	fma.rn.ftz.f32 	%f212, %f211, %f393, %f389;
+	fma.rn.ftz.f32 	%f213, %f212, %f393, %f388;
+	fma.rn.ftz.f32 	%f214, %f213, %f393, %f387;
+	fma.rn.ftz.f32 	%f215, %f214, %f393, %f386;
+	neg.f32 	%f216, %f393;
+	selp.f32 	%f217, %f216, %f12, %p11;
+	fma.rn.ftz.f32 	%f394, %f215, %f217, %f217;
+	@%p12 bra 	$L__BB0_16;
+	ex2.approx.ftz.f32 	%f218, %f394;
+	sub.f32 	%f220, %f358, %f218;
+	mov.b32 	%r31, %f220;
+	mov.b32 	%r32, %f12;
+	and.b32  	%r33, %r32, -2147483648;
+	or.b32  	%r34, %r33, %r31;
+	mov.b32 	%f394, %r34;
+$L__BB0_16:
+	.loc	1 0 0
+	cvt.u16.u32 	%rs7, %r5;
+	mov.b32 	%f6, %r11;
+	.loc	1 30 23
+	abs.ftz.f32 	%f69, %f13;
+	setp.ge.f32 	%p14, %f69, 0f3F8060FE;
+	mov.f32 	%f401, 0f3789CA3C;
+	mov.f32 	%f400, 0fB9F560B9;
+	mov.f32 	%f399, 0f3BAC840B;
+	mov.f32 	%f398, 0fBD0C8162;
+	mov.f32 	%f397, 0f3E1CF906;
+	mov.f32 	%f396, 0f3F6A937E;
+	mov.f32 	%f395, 0f3F20D842;
+	mov.f32 	%f402, %f69;
+	@%p14 bra 	$L__BB0_18;
+	mul.f32 	%f402, %f13, %f13;
+	mov.f32 	%f401, 0f38B1E96A;
+	mov.f32 	%f400, 0fBA574D20;
+	mov.f32 	%f399, 0f3BAAD5EA;
+	mov.f32 	%f398, 0fBCDC1BE7;
+	mov.f32 	%f397, 0f3DE718AF;
+	mov.f32 	%f396, 0fBEC093AC;
+	mov.f32 	%f395, 0f3E0375D3;
+$L__BB0_18:
+	.loc	1 0 0
+	cvt.f32.bf16 %r12, %rs7;
+	mul.f32 	%f14, %f6, 0f3F3504F3;
+	.loc	1 30 23
+	setp.ltu.f32 	%p15, %f69, 0f3F8060FE;
+	fma.rn.ftz.f32 	%f235, %f401, %f402, %f400;
+	fma.rn.ftz.f32 	%f236, %f235, %f402, %f399;
+	fma.rn.ftz.f32 	%f237, %f236, %f402, %f398;
+	fma.rn.ftz.f32 	%f238, %f237, %f402, %f397;
+	fma.rn.ftz.f32 	%f239, %f238, %f402, %f396;
+	fma.rn.ftz.f32 	%f240, %f239, %f402, %f395;
+	neg.f32 	%f241, %f402;
+	selp.f32 	%f242, %f241, %f13, %p14;
+	fma.rn.ftz.f32 	%f403, %f240, %f242, %f242;
+	@%p15 bra 	$L__BB0_20;
+	ex2.approx.ftz.f32 	%f243, %f403;
+	sub.f32 	%f245, %f358, %f243;
+	mov.b32 	%r35, %f245;
+	mov.b32 	%r36, %f13;
+	and.b32  	%r37, %r36, -2147483648;
+	or.b32  	%r38, %r37, %r35;
+	mov.b32 	%f403, %r38;
+$L__BB0_20:
+	.loc	1 0 0
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs8}, %r5; }
+	mov.b32 	%f7, %r12;
+	.loc	1 30 23
+	abs.ftz.f32 	%f82, %f14;
+	setp.ge.f32 	%p17, %f82, 0f3F8060FE;
+	mov.f32 	%f410, 0f3789CA3C;
+	mov.f32 	%f409, 0fB9F560B9;
+	mov.f32 	%f408, 0f3BAC840B;
+	mov.f32 	%f407, 0fBD0C8162;
+	mov.f32 	%f406, 0f3E1CF906;
+	mov.f32 	%f405, 0f3F6A937E;
+	mov.f32 	%f404, 0f3F20D842;
+	mov.f32 	%f411, %f82;
+	@%p17 bra 	$L__BB0_22;
+	mul.f32 	%f411, %f14, %f14;
+	mov.f32 	%f410, 0f38B1E96A;
+	mov.f32 	%f409, 0fBA574D20;
+	mov.f32 	%f408, 0f3BAAD5EA;
+	mov.f32 	%f407, 0fBCDC1BE7;
+	mov.f32 	%f406, 0f3DE718AF;
+	mov.f32 	%f405, 0fBEC093AC;
+	mov.f32 	%f404, 0f3E0375D3;
+$L__BB0_22:
+	.loc	1 0 0
+	cvt.f32.bf16 %r13, %rs8;
+	mul.f32 	%f15, %f7, 0f3F3504F3;
+	.loc	1 30 23
+	setp.ltu.f32 	%p18, %f82, 0f3F8060FE;
+	fma.rn.ftz.f32 	%f260, %f410, %f411, %f409;
+	fma.rn.ftz.f32 	%f261, %f260, %f411, %f408;
+	fma.rn.ftz.f32 	%f262, %f261, %f411, %f407;
+	fma.rn.ftz.f32 	%f263, %f262, %f411, %f406;
+	fma.rn.ftz.f32 	%f264, %f263, %f411, %f405;
+	fma.rn.ftz.f32 	%f265, %f264, %f411, %f404;
+	neg.f32 	%f266, %f411;
+	selp.f32 	%f267, %f266, %f14, %p17;
+	fma.rn.ftz.f32 	%f412, %f265, %f267, %f267;
+	@%p18 bra 	$L__BB0_24;
+	ex2.approx.ftz.f32 	%f268, %f412;
+	sub.f32 	%f270, %f358, %f268;
+	mov.b32 	%r39, %f270;
+	mov.b32 	%r40, %f14;
+	and.b32  	%r41, %r40, -2147483648;
+	or.b32  	%r42, %r41, %r39;
+	mov.b32 	%f412, %r42;
+$L__BB0_24:
+	.loc	1 0 0
+	mov.b32 	%f8, %r13;
+	.loc	1 30 23
+	abs.ftz.f32 	%f95, %f15;
+	setp.ge.f32 	%p20, %f95, 0f3F8060FE;
+	mov.f32 	%f419, 0f3789CA3C;
+	mov.f32 	%f418, 0fB9F560B9;
+	mov.f32 	%f417, 0f3BAC840B;
+	mov.f32 	%f416, 0fBD0C8162;
+	mov.f32 	%f415, 0f3E1CF906;
+	mov.f32 	%f414, 0f3F6A937E;
+	mov.f32 	%f413, 0f3F20D842;
+	mov.f32 	%f420, %f95;
+	@%p20 bra 	$L__BB0_26;
+	mul.f32 	%f420, %f15, %f15;
+	mov.f32 	%f419, 0f38B1E96A;
+	mov.f32 	%f418, 0fBA574D20;
+	mov.f32 	%f417, 0f3BAAD5EA;
+	mov.f32 	%f416, 0fBCDC1BE7;
+	mov.f32 	%f415, 0f3DE718AF;
+	mov.f32 	%f414, 0fBEC093AC;
+	mov.f32 	%f413, 0f3E0375D3;
+$L__BB0_26:
+	.loc	1 0 0
+	mul.f32 	%f16, %f8, 0f3F3504F3;
+	.loc	1 30 23
+	setp.ltu.f32 	%p21, %f95, 0f3F8060FE;
+	fma.rn.ftz.f32 	%f285, %f419, %f420, %f418;
+	fma.rn.ftz.f32 	%f286, %f285, %f420, %f417;
+	fma.rn.ftz.f32 	%f287, %f286, %f420, %f416;
+	fma.rn.ftz.f32 	%f288, %f287, %f420, %f415;
+	fma.rn.ftz.f32 	%f289, %f288, %f420, %f414;
+	fma.rn.ftz.f32 	%f290, %f289, %f420, %f413;
+	neg.f32 	%f291, %f420;
+	selp.f32 	%f292, %f291, %f15, %p20;
+	fma.rn.ftz.f32 	%f421, %f290, %f292, %f292;
+	@%p21 bra 	$L__BB0_28;
+	ex2.approx.ftz.f32 	%f293, %f421;
+	sub.f32 	%f295, %f358, %f293;
+	mov.b32 	%r43, %f295;
+	mov.b32 	%r44, %f15;
+	and.b32  	%r45, %r44, -2147483648;
+	or.b32  	%r46, %r45, %r43;
+	mov.b32 	%f421, %r46;
+$L__BB0_28:
+	abs.ftz.f32 	%f108, %f16;
+	setp.ge.f32 	%p23, %f108, 0f3F8060FE;
+	mov.f32 	%f428, 0f3789CA3C;
+	mov.f32 	%f427, 0fB9F560B9;
+	mov.f32 	%f426, 0f3BAC840B;
+	mov.f32 	%f425, 0fBD0C8162;
+	mov.f32 	%f424, 0f3E1CF906;
+	mov.f32 	%f423, 0f3F6A937E;
+	mov.f32 	%f422, 0f3F20D842;
+	mov.f32 	%f429, %f108;
+	@%p23 bra 	$L__BB0_30;
+	mul.f32 	%f429, %f16, %f16;
+	mov.f32 	%f428, 0f38B1E96A;
+	mov.f32 	%f427, 0fBA574D20;
+	mov.f32 	%f426, 0f3BAAD5EA;
+	mov.f32 	%f425, 0fBCDC1BE7;
+	mov.f32 	%f424, 0f3DE718AF;
+	mov.f32 	%f423, 0fBEC093AC;
+	mov.f32 	%f422, 0f3E0375D3;
+$L__BB0_30:
+	.loc	1 0 23
+	ld.param.u64 	%rd2, [triton__0d1d2de_param_1];
+	cvt.s64.s32 	%rd1, %r18;
+	.loc	1 30 23
+	setp.ltu.f32 	%p24, %f108, 0f3F8060FE;
+	fma.rn.ftz.f32 	%f310, %f428, %f429, %f427;
+	fma.rn.ftz.f32 	%f311, %f310, %f429, %f426;
+	fma.rn.ftz.f32 	%f312, %f311, %f429, %f425;
+	fma.rn.ftz.f32 	%f313, %f312, %f429, %f424;
+	fma.rn.ftz.f32 	%f314, %f313, %f429, %f423;
+	fma.rn.ftz.f32 	%f315, %f314, %f429, %f422;
+	neg.f32 	%f316, %f429;
+	selp.f32 	%f317, %f316, %f16, %p23;
+	fma.rn.ftz.f32 	%f430, %f315, %f317, %f317;
+	@%p24 bra 	$L__BB0_32;
+	ex2.approx.ftz.f32 	%f318, %f430;
+	sub.f32 	%f320, %f358, %f318;
+	mov.b32 	%r47, %f320;
+	mov.b32 	%r48, %f16;
+	and.b32  	%r49, %r48, -2147483648;
+	or.b32  	%r50, %r49, %r47;
+	mov.b32 	%f430, %r50;
+$L__BB0_32:
+	.loc	1 27 18
+	mul.f32 	%f321, %f8, 0f3F000000;
+	mul.f32 	%f322, %f7, 0f3F000000;
+	mul.f32 	%f323, %f6, 0f3F000000;
+	mul.f32 	%f324, %f5, 0f3F000000;
+	mul.f32 	%f325, %f4, 0f3F000000;
+	mul.f32 	%f326, %f3, 0f3F000000;
+	mul.f32 	%f327, %f2, 0f3F000000;
+	mul.f32 	%f328, %f1, 0f3F000000;
+	.loc	1 32 18
+	add.f32 	%f329, %f367, 0f3F800000;
+	add.f32 	%f330, %f376, 0f3F800000;
+	add.f32 	%f331, %f385, 0f3F800000;
+	add.f32 	%f332, %f394, 0f3F800000;
+	add.f32 	%f333, %f403, 0f3F800000;
+	add.f32 	%f334, %f412, 0f3F800000;
+	add.f32 	%f335, %f421, 0f3F800000;
+	add.f32 	%f336, %f430, 0f3F800000;
+	.loc	1 33 18
+	mul.f32 	%f337, %f328, %f329;
+	mul.f32 	%f338, %f327, %f330;
+	mul.f32 	%f339, %f326, %f331;
+	mul.f32 	%f340, %f325, %f332;
+	mul.f32 	%f341, %f324, %f333;
+	mul.f32 	%f342, %f323, %f334;
+	mul.f32 	%f343, %f322, %f335;
+	mul.f32 	%f344, %f321, %f336;
+	.loc	1 35 25
+	shl.b64 	%rd7, %rd1, 1;
+	add.s64 	%rd6, %rd2, %rd7;
+	.loc	1 35 37
+	mov.b32 	%r51, %f337;
+	cvt.rn.bf16.f32 %rs9, %r51;
+	mov.b32 	%r52, %f338;
+	cvt.rn.bf16.f32 %rs10, %r52;
+	mov.b32 	%r53, %f339;
+	cvt.rn.bf16.f32 %rs11, %r53;
+	mov.b32 	%r54, %f340;
+	cvt.rn.bf16.f32 %rs12, %r54;
+	mov.b32 	%r55, %f341;
+	cvt.rn.bf16.f32 %rs13, %r55;
+	mov.b32 	%r56, %f342;
+	cvt.rn.bf16.f32 %rs14, %r56;
+	mov.b32 	%r57, %f343;
+	cvt.rn.bf16.f32 %rs15, %r57;
+	mov.b32 	%r58, %f344;
+	cvt.rn.bf16.f32 %rs16, %r58;
+	mov.b32 	%r63, {%rs9, %rs10};
+	mov.b32 	%r64, {%rs11, %rs12};
+	mov.b32 	%r65, {%rs13, %rs14};
+	mov.b32 	%r66, {%rs15, %rs16};
+	@%p1 st.global.v4.b32 [ %rd6 + 0 ], { %r63, %r64, %r65, %r66 };
+	.loc	1 35 4
+	ret;
+$L__tmp1:
+$L__func_end0:
+
+}
+	// .globl	__nv_erff
+.visible .func  (.param .b32 func_retval0) __nv_erff(
+	.param .b32 __nv_erff_param_0
+)
+{
+	.reg .pred 	%p<4>;
+	.reg .b32 	%r<5>;
+	.reg .f32 	%f<49>;
+$L__func_begin1:
+
+	ld.param.f32 	%f14, [__nv_erff_param_0];
+	abs.ftz.f32 	%f1, %f14;
+	setp.ge.f32 	%p1, %f1, 0f3F8060FE;
+	mov.f32 	%f46, 0f3789CA3C;
+	mov.f32 	%f45, 0fB9F560B9;
+	mov.f32 	%f44, 0f3BAC840B;
+	mov.f32 	%f43, 0fBD0C8162;
+	mov.f32 	%f42, 0f3E1CF906;
+	mov.f32 	%f41, 0f3F6A937E;
+	mov.f32 	%f40, 0f3F20D842;
+	mov.f32 	%f47, %f1;
+	@%p1 bra 	$L__BB1_2;
+	mul.f32 	%f47, %f14, %f14;
+	mov.f32 	%f46, 0f38B1E96A;
+	mov.f32 	%f45, 0fBA574D20;
+	mov.f32 	%f44, 0f3BAAD5EA;
+	mov.f32 	%f43, 0fBCDC1BE7;
+	mov.f32 	%f42, 0f3DE718AF;
+	mov.f32 	%f41, 0fBEC093AC;
+	mov.f32 	%f40, 0f3E0375D3;
+$L__BB1_2:
+	setp.ltu.f32 	%p2, %f1, 0f3F8060FE;
+	fma.rn.ftz.f32 	%f29, %f46, %f47, %f45;
+	fma.rn.ftz.f32 	%f30, %f29, %f47, %f44;
+	fma.rn.ftz.f32 	%f31, %f30, %f47, %f43;
+	fma.rn.ftz.f32 	%f32, %f31, %f47, %f42;
+	fma.rn.ftz.f32 	%f33, %f32, %f47, %f41;
+	fma.rn.ftz.f32 	%f34, %f33, %f47, %f40;
+	neg.f32 	%f35, %f47;
+	selp.f32 	%f36, %f35, %f14, %p1;
+	fma.rn.ftz.f32 	%f48, %f34, %f36, %f36;
+	@%p2 bra 	$L__BB1_4;
+	ex2.approx.ftz.f32 	%f37, %f48;
+	mov.f32 	%f38, 0f3F800000;
+	sub.f32 	%f39, %f38, %f37;
+	mov.b32 	%r1, %f39;
+	mov.b32 	%r2, %f14;
+	and.b32  	%r3, %r2, -2147483648;
+	or.b32  	%r4, %r3, %r1;
+	mov.b32 	%f48, %r4;
+$L__BB1_4:
+	st.param.f32 	[func_retval0+0], %f48;
+	ret;
+$L__func_end1:
+
+}
+	.file	1 "/tmp/torchinductor_root/jf/cjfoqo3nutni5cmtw4brla34cz45fusadehkxfkr2fie2qgo7vwt.py"
+	.section	.debug_abbrev
+	{
+.b8 1
+.b8 17
+.b8 1
+.b8 37
+.b8 8
+.b8 19
+.b8 5
+.b8 3
+.b8 8
+.b8 16
+.b8 6
+.b8 27
+.b8 8
+.b8 180
+.b8 66
+.b8 12
+.b8 17
+.b8 1
+.b8 18
+.b8 1
+.b8 0
+.b8 0
+.b8 2
+.b8 46
+.b8 0
+.b8 17
+.b8 1
+.b8 18
+.b8 1
+.b8 64
+.b8 10
+.b8 135
+.b8 64
+.b8 8
+.b8 3
+.b8 8
+.b8 58
+.b8 11
+.b8 59
+.b8 11
+.b8 63
+.b8 12
+.b8 0
+.b8 0
+.b8 0
+	}
+	.section	.debug_info
+	{
+.b32 176
+.b8 2
+.b8 0
+.b32 .debug_abbrev
+.b8 8
+.b8 1
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 0
+.b8 2
+.b8 0
+.b8 99
+.b8 106
+.b8 102
+.b8 111
+.b8 113
+.b8 111
+.b8 51
+.b8 110
+.b8 117
+.b8 116
+.b8 110
+.b8 105
+.b8 53
+.b8 99
+.b8 109
+.b8 116
+.b8 119
+.b8 52
+.b8 98
+.b8 114
+.b8 108
+.b8 97
+.b8 51
+.b8 52
+.b8 99
+.b8 122
+.b8 52
+.b8 53
+.b8 102
+.b8 117
+.b8 115
+.b8 97
+.b8 100
+.b8 101
+.b8 104
+.b8 107
+.b8 120
+.b8 102
+.b8 107
+.b8 114
+.b8 50
+.b8 102
+.b8 105
+.b8 101
+.b8 50
+.b8 113
+.b8 103
+.b8 111
+.b8 55
+.b8 118
+.b8 119
+.b8 116
+.b8 46
+.b8 112
+.b8 121
+.b8 0
+.b32 .debug_line
+.b8 47
+.b8 116
+.b8 109
+.b8 112
+.b8 47
+.b8 116
+.b8 111
+.b8 114
+.b8 99
+.b8 104
+.b8 105
+.b8 110
+.b8 100
+.b8 117
+.b8 99
+.b8 116
+.b8 111
+.b8 114
+.b8 95
+.b8 114
+.b8 111
+.b8 111
+.b8 116
+.b8 47
+.b8 106
+.b8 102
+.b8 0
+.b8 1
+.b64 $L__func_begin0
+.b64 $L__func_end0
+.b8 2
+.b64 $L__func_begin0
+.b64 $L__func_end0
+.b8 1
+.b8 156
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 95
+.b8 95
+.b8 48
+.b8 100
+.b8 49
+.b8 100
+.b8 50
+.b8 100
+.b8 101
+.b8 0
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 95
+.b8 95
+.b8 48
+.b8 100
+.b8 49
+.b8 100
+.b8 50
+.b8 100
+.b8 101
+.b8 0
+.b8 1
+.b8 18
+.b8 1
+.b8 0
+	}
+	.section	.debug_pubnames
+	{
+.b32 $L__pubNames_end0-$L__pubNames_start0
+$L__pubNames_start0:
+.b8 2
+.b8 0
+.b32 .debug_info
+.b32 180
+.b32 125
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 95
+.b8 95
+.b8 48
+.b8 100
+.b8 49
+.b8 100
+.b8 50
+.b8 100
+.b8 101
+.b8 0
+.b32 0
+$L__pubNames_end0:
+	}
+	.section	.debug_pubtypes
+	{
+.b32 $L__pubTypes_end0-$L__pubTypes_start0
+$L__pubTypes_start0:
+.b8 2
+.b8 0
+.b32 .debug_info
+.b32 180
+.b32 0
+$L__pubTypes_end0:
+	}
+	.section	.debug_loc	{	}

.triton/dump/415aac87553b7d064f52694fa7254686/triton_.ttir

@@ -0,0 +1,27 @@
+module {
+  tt.func public @triton__0d1d2de(%arg0: !tt.ptr<bf16, 1> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<bf16, 1> {tt.divisibility = 16 : i32}, %arg2: i32 {tt.divisibility = 16 : i32, tt.max_divisibility = 16 : i32}) attributes {noinline = false} {
+    %cst = arith.constant dense<1.000000e+00> : tensor<1024xf32>
+    %cst_0 = arith.constant dense<0.707106769> : tensor<1024xf32>
+    %cst_1 = arith.constant dense<5.000000e-01> : tensor<1024xf32>
+    %c1024_i32 = arith.constant 1024 : i32
+    %0 = tt.get_program_id x : i32
+    %1 = arith.muli %0, %c1024_i32 : i32
+    %2 = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32>
+    %3 = tt.splat %1 : (i32) -> tensor<1024xi32>
+    %4 = arith.addi %3, %2 : tensor<1024xi32>
+    %5 = tt.splat %arg0 : (!tt.ptr<bf16, 1>) -> tensor<1024x!tt.ptr<bf16, 1>>
+    %6 = tt.addptr %5, %4 : tensor<1024x!tt.ptr<bf16, 1>>, tensor<1024xi32>
+    %7 = tt.load %6 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<1024xbf16>
+    %8 = arith.extf %7 : tensor<1024xbf16> to tensor<1024xf32>
+    %9 = arith.mulf %8, %cst_1 : tensor<1024xf32>
+    %10 = arith.mulf %8, %cst_0 : tensor<1024xf32>
+    %11 = tt.extern_elementwise %10 {libname = "libdevice", libpath = "/usr/local/lib/python3.10/dist-packages/triton/language/../third_party/cuda/lib/libdevice.10.bc", pure = true, symbol = "__nv_erff"} : (tensor<1024xf32>) -> tensor<1024xf32>
+    %12 = arith.addf %11, %cst : tensor<1024xf32>
+    %13 = arith.mulf %9, %12 : tensor<1024xf32>
+    %14 = tt.splat %arg1 : (!tt.ptr<bf16, 1>) -> tensor<1024x!tt.ptr<bf16, 1>>
+    %15 = tt.addptr %14, %4 : tensor<1024x!tt.ptr<bf16, 1>>, tensor<1024xi32>
+    %16 = arith.truncf %13 : tensor<1024xf32> to tensor<1024xbf16>
+    tt.store %15, %16 {cache = 1 : i32, evict = 1 : i32} : tensor<1024xbf16>
+    tt.return
+  }
+}

.triton/dump/645565eaba0a18dd23ef200fe9abb0c0/triton_.ttir

@@ -0,0 +1,89 @@
+module {
+  tt.func public @triton__0d1d2d3d4d5d6d7d8de9de(%arg0: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<i64, 1> {tt.divisibility = 16 : i32}, %arg2: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg3: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg4: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg5: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg6: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg7: !tt.ptr<bf16, 1> {tt.divisibility = 16 : i32}, %arg8: i32 {tt.divisibility = 16 : i32, tt.max_divisibility = 16 : i32}, %arg9: i32 {tt.divisibility = 16 : i32, tt.max_divisibility = 16 : i32}) attributes {noinline = false} {
+    %c512_i32 = arith.constant 512 : i32
+    %c256_i32 = arith.constant 256 : i32
+    %cst = arith.constant 0.000000e+00 : f32
+    %cst_0 = arith.constant 2.560000e+02 : f32
+    %cst_1 = arith.constant 9.99999974E-6 : f32
+    %cst_2 = arith.constant dense<0.000000e+00> : tensor<256xf32>
+    %cst_3 = arith.constant dense<256> : tensor<1xi64>
+    %cst_4 = arith.constant dense<50257> : tensor<1xi64>
+    %cst_5 = arith.constant dense<0> : tensor<1xi64>
+    %cst_6 = arith.constant dense<256> : tensor<256xi32>
+    %0 = tt.get_program_id x : i32
+    %1 = tt.make_range {end = 256 : i32, start = 0 : i32} : tensor<256xi32>
+    %2 = arith.cmpi slt, %1, %cst_6 : tensor<256xi32>
+    %3 = arith.remsi %0, %c512_i32 : i32
+    %4 = tt.addptr %arg1, %0 : !tt.ptr<i64, 1>, i32
+    %5 = tt.splat %4 : (!tt.ptr<i64, 1>) -> tensor<1x!tt.ptr<i64, 1>>
+    %6 = tt.load %5 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<1xi64>
+    %7 = arith.muli %3, %c256_i32 : i32
+    %8 = tt.splat %7 : (i32) -> tensor<256xi32>
+    %9 = arith.addi %1, %8 : tensor<256xi32>
+    %10 = tt.splat %arg3 : (!tt.ptr<f32, 1>) -> tensor<256x!tt.ptr<f32, 1>>
+    %11 = tt.addptr %10, %9 : tensor<256x!tt.ptr<f32, 1>>, tensor<256xi32>
+    %12 = tt.load %11, %2, %cst_2 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<256xf32>
+    %13 = tt.splat %arg4 : (!tt.ptr<f32, 1>) -> tensor<256x!tt.ptr<f32, 1>>
+    %14 = tt.addptr %13, %1 : tensor<256x!tt.ptr<f32, 1>>, tensor<256xi32>
+    %15 = tt.load %14, %2, %cst_2 {cache = 1 : i32, evict = 3 : i32, isVolatile = false} : tensor<256xf32>
+    %16 = arith.addi %6, %cst_4 : tensor<1xi64>
+    %17 = arith.cmpi slt, %6, %cst_5 : tensor<1xi64>
+    %18 = arith.select %17, %16, %6 : tensor<1xi1>, tensor<1xi64>
+    %19 = arith.cmpi sge, %18, %cst_5 : tensor<1xi64>
+    %20 = arith.cmpi slt, %18, %cst_4 : tensor<1xi64>
+    %21 = arith.andi %19, %20 : tensor<1xi1>
+    tt.assert %21, "index out of bounds: 0 <= tmp3 < 50257", "<frozen importlib._bootstrap_external>", "_call_with_frames_removed", 883 : tensor<1xi1>
+    %22 = arith.muli %18, %cst_3 : tensor<1xi64>
+    %23 = tt.broadcast %22 : (tensor<1xi64>) -> tensor<256xi64>
+    %24 = arith.extsi %1 : tensor<256xi32> to tensor<256xi64>
+    %25 = arith.addi %24, %23 : tensor<256xi64>
+    %26 = tt.splat %arg2 : (!tt.ptr<f32, 1>) -> tensor<256x!tt.ptr<f32, 1>>
+    %27 = tt.addptr %26, %25 : tensor<256x!tt.ptr<f32, 1>>, tensor<256xi64>
+    %28 = tt.load %27, %2, %cst_2 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<256xf32>
+    %29 = arith.addf %28, %12 : tensor<256xf32>
+    %30 = arith.select %2, %29, %cst_2 : tensor<256xi1>, tensor<256xf32>
+    %31 = "tt.reduce"(%30) <{axis = 0 : i32}> ({
+    ^bb0(%arg10: f32, %arg11: f32):
+      %60 = arith.addf %arg10, %arg11 : f32
+      tt.reduce.return %60 : f32
+    }) : (tensor<256xf32>) -> f32
+    %32 = arith.addf %31, %cst : f32
+    %33 = arith.divf %32, %cst_0 : f32
+    %34 = tt.splat %33 : (f32) -> tensor<1xf32>
+    %35 = tt.splat %33 : (f32) -> tensor<256xf32>
+    %36 = arith.subf %29, %35 : tensor<256xf32>
+    %37 = arith.mulf %36, %36 : tensor<256xf32>
+    %38 = arith.select %2, %37, %cst_2 : tensor<256xi1>, tensor<256xf32>
+    %39 = "tt.reduce"(%38) <{axis = 0 : i32}> ({
+    ^bb0(%arg10: f32, %arg11: f32):
+      %60 = arith.addf %arg10, %arg11 : f32
+      tt.reduce.return %60 : f32
+    }) : (tensor<256xf32>) -> f32
+    %40 = arith.addf %39, %cst : f32
+    %41 = arith.divf %40, %cst_0 : f32
+    %42 = arith.addf %41, %cst_1 : f32
+    %43 = tt.extern_elementwise %42 {libname = "libdevice", libpath = "/usr/local/lib/python3.10/dist-packages/triton/language/../third_party/cuda/lib/libdevice.10.bc", pure = true, symbol = "__nv_rsqrtf"} : (f32) -> f32
+    %44 = tt.splat %43 : (f32) -> tensor<1xf32>
+    %45 = tt.splat %43 : (f32) -> tensor<256xf32>
+    %46 = arith.mulf %36, %45 : tensor<256xf32>
+    %47 = arith.mulf %46, %15 : tensor<256xf32>
+    %48 = arith.muli %0, %c256_i32 : i32
+    %49 = tt.splat %48 : (i32) -> tensor<256xi32>
+    %50 = arith.addi %1, %49 : tensor<256xi32>
+    %51 = tt.splat %arg5 : (!tt.ptr<f32, 1>) -> tensor<256x!tt.ptr<f32, 1>>
+    %52 = tt.addptr %51, %50 : tensor<256x!tt.ptr<f32, 1>>, tensor<256xi32>
+    tt.store %52, %29, %2 {cache = 1 : i32, evict = 1 : i32} : tensor<256xf32>
+    gpu.barrier
+    %53 = tt.addptr %arg0, %0 : !tt.ptr<f32, 1>, i32
+    %54 = tt.splat %53 : (!tt.ptr<f32, 1>) -> tensor<1x!tt.ptr<f32, 1>>
+    tt.store %54, %44 {cache = 1 : i32, evict = 1 : i32} : tensor<1xf32>
+    %55 = tt.splat %arg7 : (!tt.ptr<bf16, 1>) -> tensor<256x!tt.ptr<bf16, 1>>
+    %56 = tt.addptr %55, %50 : tensor<256x!tt.ptr<bf16, 1>>, tensor<256xi32>
+    %57 = arith.truncf %47 : tensor<256xf32> to tensor<256xbf16>
+    tt.store %56, %57, %2 {cache = 1 : i32, evict = 1 : i32} : tensor<256xbf16>
+    %58 = tt.addptr %arg6, %0 : !tt.ptr<f32, 1>, i32
+    %59 = tt.splat %58 : (!tt.ptr<f32, 1>) -> tensor<1x!tt.ptr<f32, 1>>
+    tt.store %59, %34 {cache = 1 : i32, evict = 1 : i32} : tensor<1xf32>
+    tt.return
+  }
+}

.triton/dump/7dc5bb3e5c2bb99527fff34c6fba7810/triton_.ttgir

@@ -0,0 +1,18 @@
+#blocked = #triton_gpu.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0], CTAsPerCGA = [1], CTASplitNum = [1], CTAOrder = [0]}>
+module attributes {"triton_gpu.compute-capability" = 89 : i32, "triton_gpu.num-ctas" = 1 : i32, "triton_gpu.num-warps" = 4 : i32, "triton_gpu.threads-per-warp" = 32 : i32} {
+  tt.func public @triton__0d1de(%arg0: !tt.ptr<i64, 1> {tt.divisibility = 16 : i32}, %arg1: i32 {tt.divisibility = 16 : i32, tt.max_divisibility = 16 : i32}) attributes {noinline = false} {
+    %cst = arith.constant dense<512> : tensor<128xi32, #blocked>
+    %c128_i32 = arith.constant 128 : i32
+    %0 = tt.get_program_id x : i32
+    %1 = arith.muli %0, %c128_i32 : i32
+    %2 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32, #blocked>
+    %3 = tt.splat %1 : (i32) -> tensor<128xi32, #blocked>
+    %4 = arith.addi %3, %2 : tensor<128xi32, #blocked>
+    %5 = arith.cmpi slt, %4, %cst : tensor<128xi32, #blocked>
+    %6 = tt.splat %arg0 : (!tt.ptr<i64, 1>) -> tensor<128x!tt.ptr<i64, 1>, #blocked>
+    %7 = tt.addptr %6, %4 : tensor<128x!tt.ptr<i64, 1>, #blocked>, tensor<128xi32, #blocked>
+    %8 = arith.extsi %4 : tensor<128xi32, #blocked> to tensor<128xi64, #blocked>
+    tt.store %7, %8, %5 {cache = 1 : i32, evict = 1 : i32} : tensor<128xi64, #blocked>
+    tt.return
+  }
+}

.triton/dump/8c4bac4d904709a8b7e8c698132d974c/triton_.ttir

@@ -0,0 +1,17 @@
+module {
+  tt.func public @triton__0d1de(%arg0: !tt.ptr<i64, 1> {tt.divisibility = 16 : i32}, %arg1: i32 {tt.divisibility = 16 : i32, tt.max_divisibility = 16 : i32}) attributes {noinline = false} {
+    %cst = arith.constant dense<512> : tensor<256xi32>
+    %c256_i32 = arith.constant 256 : i32
+    %0 = tt.get_program_id x : i32
+    %1 = arith.muli %0, %c256_i32 : i32
+    %2 = tt.make_range {end = 256 : i32, start = 0 : i32} : tensor<256xi32>
+    %3 = tt.splat %1 : (i32) -> tensor<256xi32>
+    %4 = arith.addi %3, %2 : tensor<256xi32>
+    %5 = arith.cmpi slt, %4, %cst : tensor<256xi32>
+    %6 = tt.splat %arg0 : (!tt.ptr<i64, 1>) -> tensor<256x!tt.ptr<i64, 1>>
+    %7 = tt.addptr %6, %4 : tensor<256x!tt.ptr<i64, 1>>, tensor<256xi32>
+    %8 = arith.extsi %4 : tensor<256xi32> to tensor<256xi64>
+    tt.store %7, %8, %5 {cache = 1 : i32, evict = 1 : i32} : tensor<256xi64>
+    tt.return
+  }
+}

.triton/dump/93e5abc5363b9438178c618128714f73/triton_.ptx

@@ -0,0 +1,861 @@
+//
+// Generated by LLVM NVPTX Back-End
+//
+
+.version 8.2
+.target sm_89
+.address_size 64
+
+	// .globl	triton__0d1d2de
+.global .align 1 .b8 _$_str[11] = {95, 95, 67, 85, 68, 65, 95, 70, 84, 90, 0};
+
+.visible .entry triton__0d1d2de(
+	.param .u64 triton__0d1d2de_param_0,
+	.param .u64 triton__0d1d2de_param_1,
+	.param .u32 triton__0d1d2de_param_2
+)
+.maxntid 128, 1, 1
+{
+	.reg .pred 	%p<28>;
+	.reg .b16 	%rs<25>;
+	.reg .b32 	%r<79>;
+	.reg .f32 	%f<487>;
+	.reg .b64 	%rd<8>;
+	.loc	1 18 0
+$L__func_begin0:
+	.loc	1 18 0
+
+	ld.param.u64 	%rd4, [triton__0d1d2de_param_0];
+	ld.param.u64 	%rd5, [triton__0d1d2de_param_1];
+$L__tmp0:
+	.loc	1 21 36
+	mov.u32 	%r26, %tid.x;
+	shl.b32 	%r27, %r26, 3;
+	and.b32  	%r28, %r27, 1016;
+	.loc	1 20 28
+	mov.u32 %r1, %ctaid.x;
+	.loc	1 20 33
+	shl.b32 	%r29, %r1, 10;
+	.loc	1 21 23
+	or.b32  	%r30, %r29, %r28;
+	.loc	1 24 34
+	mul.wide.s32 	%rd6, %r30, 2;
+	add.s64 	%rd7, %rd4, %rd6;
+	mov.pred 	%p1, -1;
+	.loc	1 24 39
+	mov.u32 %r2, 0x0;
+	mov.u32 %r3, 0x0;
+	mov.u32 %r4, 0x0;
+	mov.u32 %r5, 0x0;
+	@%p1 ld.global.v4.b32 { %r2, %r3, %r4, %r5 }, [ %rd7 + 0 ];
+	.loc	1 25 30
+	add.s64 	%rd3, %rd5, %rd6;
+	.loc	1 25 35
+	mov.u32 %r14, 0x0;
+	mov.u32 %r15, 0x0;
+	mov.u32 %r16, 0x0;
+	mov.u32 %r17, 0x0;
+	@%p1 ld.global.v4.b32 { %r14, %r15, %r16, %r17 }, [ %rd3 + 0 ];
+	cvt.u16.u32 	%rs9, %r14;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs10}, %r14; }
+	cvt.u16.u32 	%rs11, %r15;
+	.loc	1 25 44
+	cvt.f32.bf16 %r18, %rs9;
+	mov.b32 	%f9, %r18;
+	cvt.f32.bf16 %r19, %rs10;
+	mov.b32 	%f10, %r19;
+	.loc	1 29 18
+	mul.f32 	%f17, %f9, 0f3F3504F3;
+	.loc	1 30 23
+	abs.ftz.f32 	%f25, %f17;
+	setp.ge.f32 	%p3, %f25, 0f3F8060FE;
+	mov.f32 	%f421, 0f3789CA3C;
+	mov.f32 	%f420, 0fB9F560B9;
+	mov.f32 	%f419, 0f3BAC840B;
+	mov.f32 	%f418, 0fBD0C8162;
+	mov.f32 	%f417, 0f3E1CF906;
+	mov.f32 	%f416, 0f3F6A937E;
+	mov.f32 	%f415, 0f3F20D842;
+	mov.f32 	%f422, %f25;
+	@%p3 bra 	$L__BB0_2;
+	.loc	1 0 23
+	mov.f32 	%f421, 0f38B1E96A;
+	mov.f32 	%f420, 0fBA574D20;
+	mov.f32 	%f419, 0f3BAAD5EA;
+	mov.f32 	%f418, 0fBCDC1BE7;
+	mov.f32 	%f417, 0f3DE718AF;
+	mov.f32 	%f416, 0fBEC093AC;
+	mov.f32 	%f415, 0f3E0375D3;
+	.loc	1 30 23
+	mul.f32 	%f422, %f17, %f17;
+$L__BB0_2:
+	.loc	1 0 0
+	cvt.f32.bf16 %r20, %rs11;
+	mul.f32 	%f18, %f10, 0f3F3504F3;
+	.loc	1 30 23
+	setp.ltu.f32 	%p4, %f25, 0f3F8060FE;
+	fma.rn.ftz.f32 	%f143, %f421, %f422, %f420;
+	fma.rn.ftz.f32 	%f144, %f143, %f422, %f419;
+	fma.rn.ftz.f32 	%f145, %f144, %f422, %f418;
+	fma.rn.ftz.f32 	%f146, %f145, %f422, %f417;
+	fma.rn.ftz.f32 	%f147, %f146, %f422, %f416;
+	fma.rn.ftz.f32 	%f148, %f147, %f422, %f415;
+	neg.f32 	%f149, %f422;
+	selp.f32 	%f150, %f149, %f17, %p3;
+	fma.rn.ftz.f32 	%f423, %f148, %f150, %f150;
+	mov.f32 	%f414, 0f3F800000;
+	@%p4 bra 	$L__BB0_4;
+	ex2.approx.ftz.f32 	%f151, %f423;
+	sub.f32 	%f153, %f414, %f151;
+	mov.b32 	%r31, %f153;
+	mov.b32 	%r32, %f17;
+	and.b32  	%r33, %r32, -2147483648;
+	or.b32  	%r34, %r33, %r31;
+	mov.b32 	%f423, %r34;
+$L__BB0_4:
+	.loc	1 0 0
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs12}, %r15; }
+	mov.b32 	%f11, %r20;
+	.loc	1 30 23
+	abs.ftz.f32 	%f38, %f18;
+	setp.ge.f32 	%p6, %f38, 0f3F8060FE;
+	mov.f32 	%f430, 0f3789CA3C;
+	mov.f32 	%f429, 0fB9F560B9;
+	mov.f32 	%f428, 0f3BAC840B;
+	mov.f32 	%f427, 0fBD0C8162;
+	mov.f32 	%f426, 0f3E1CF906;
+	mov.f32 	%f425, 0f3F6A937E;
+	mov.f32 	%f424, 0f3F20D842;
+	mov.f32 	%f431, %f38;
+	@%p6 bra 	$L__BB0_6;
+	mul.f32 	%f431, %f18, %f18;
+	mov.f32 	%f430, 0f38B1E96A;
+	mov.f32 	%f429, 0fBA574D20;
+	mov.f32 	%f428, 0f3BAAD5EA;
+	mov.f32 	%f427, 0fBCDC1BE7;
+	mov.f32 	%f426, 0f3DE718AF;
+	mov.f32 	%f425, 0fBEC093AC;
+	mov.f32 	%f424, 0f3E0375D3;
+$L__BB0_6:
+	.loc	1 0 0
+	cvt.f32.bf16 %r21, %rs12;
+	mul.f32 	%f19, %f11, 0f3F3504F3;
+	.loc	1 30 23
+	setp.ltu.f32 	%p7, %f38, 0f3F8060FE;
+	fma.rn.ftz.f32 	%f168, %f430, %f431, %f429;
+	fma.rn.ftz.f32 	%f169, %f168, %f431, %f428;
+	fma.rn.ftz.f32 	%f170, %f169, %f431, %f427;
+	fma.rn.ftz.f32 	%f171, %f170, %f431, %f426;
+	fma.rn.ftz.f32 	%f172, %f171, %f431, %f425;
+	fma.rn.ftz.f32 	%f173, %f172, %f431, %f424;
+	neg.f32 	%f174, %f431;
+	selp.f32 	%f175, %f174, %f18, %p6;
+	fma.rn.ftz.f32 	%f432, %f173, %f175, %f175;
+	@%p7 bra 	$L__BB0_8;
+	ex2.approx.ftz.f32 	%f176, %f432;
+	sub.f32 	%f178, %f414, %f176;
+	mov.b32 	%r35, %f178;
+	mov.b32 	%r36, %f18;
+	and.b32  	%r37, %r36, -2147483648;
+	or.b32  	%r38, %r37, %r35;
+	mov.b32 	%f432, %r38;
+$L__BB0_8:
+	.loc	1 0 0
+	cvt.u16.u32 	%rs13, %r16;
+	mov.b32 	%f12, %r21;
+	.loc	1 30 23
+	abs.ftz.f32 	%f51, %f19;
+	setp.ge.f32 	%p9, %f51, 0f3F8060FE;
+	mov.f32 	%f439, 0f3789CA3C;
+	mov.f32 	%f438, 0fB9F560B9;
+	mov.f32 	%f437, 0f3BAC840B;
+	mov.f32 	%f436, 0fBD0C8162;
+	mov.f32 	%f435, 0f3E1CF906;
+	mov.f32 	%f434, 0f3F6A937E;
+	mov.f32 	%f433, 0f3F20D842;
+	mov.f32 	%f440, %f51;
+	@%p9 bra 	$L__BB0_10;
+	mul.f32 	%f440, %f19, %f19;
+	mov.f32 	%f439, 0f38B1E96A;
+	mov.f32 	%f438, 0fBA574D20;
+	mov.f32 	%f437, 0f3BAAD5EA;
+	mov.f32 	%f436, 0fBCDC1BE7;
+	mov.f32 	%f435, 0f3DE718AF;
+	mov.f32 	%f434, 0fBEC093AC;
+	mov.f32 	%f433, 0f3E0375D3;
+$L__BB0_10:
+	.loc	1 0 0
+	cvt.f32.bf16 %r22, %rs13;
+	mul.f32 	%f20, %f12, 0f3F3504F3;
+	.loc	1 30 23
+	setp.ltu.f32 	%p10, %f51, 0f3F8060FE;
+	fma.rn.ftz.f32 	%f193, %f439, %f440, %f438;
+	fma.rn.ftz.f32 	%f194, %f193, %f440, %f437;
+	fma.rn.ftz.f32 	%f195, %f194, %f440, %f436;
+	fma.rn.ftz.f32 	%f196, %f195, %f440, %f435;
+	fma.rn.ftz.f32 	%f197, %f196, %f440, %f434;
+	fma.rn.ftz.f32 	%f198, %f197, %f440, %f433;
+	neg.f32 	%f199, %f440;
+	selp.f32 	%f200, %f199, %f19, %p9;
+	fma.rn.ftz.f32 	%f441, %f198, %f200, %f200;
+	@%p10 bra 	$L__BB0_12;
+	ex2.approx.ftz.f32 	%f201, %f441;
+	sub.f32 	%f203, %f414, %f201;
+	mov.b32 	%r39, %f203;
+	mov.b32 	%r40, %f19;
+	and.b32  	%r41, %r40, -2147483648;
+	or.b32  	%r42, %r41, %r39;
+	mov.b32 	%f441, %r42;
+$L__BB0_12:
+	.loc	1 0 0
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs14}, %r16; }
+	mov.b32 	%f13, %r22;
+	.loc	1 30 23
+	abs.ftz.f32 	%f64, %f20;
+	setp.ge.f32 	%p12, %f64, 0f3F8060FE;
+	mov.f32 	%f448, 0f3789CA3C;
+	mov.f32 	%f447, 0fB9F560B9;
+	mov.f32 	%f446, 0f3BAC840B;
+	mov.f32 	%f445, 0fBD0C8162;
+	mov.f32 	%f444, 0f3E1CF906;
+	mov.f32 	%f443, 0f3F6A937E;
+	mov.f32 	%f442, 0f3F20D842;
+	mov.f32 	%f449, %f64;
+	@%p12 bra 	$L__BB0_14;
+	mul.f32 	%f449, %f20, %f20;
+	mov.f32 	%f448, 0f38B1E96A;
+	mov.f32 	%f447, 0fBA574D20;
+	mov.f32 	%f446, 0f3BAAD5EA;
+	mov.f32 	%f445, 0fBCDC1BE7;
+	mov.f32 	%f444, 0f3DE718AF;
+	mov.f32 	%f443, 0fBEC093AC;
+	mov.f32 	%f442, 0f3E0375D3;
+$L__BB0_14:
+	.loc	1 0 0
+	cvt.f32.bf16 %r23, %rs14;
+	mul.f32 	%f21, %f13, 0f3F3504F3;
+	.loc	1 30 23
+	setp.ltu.f32 	%p13, %f64, 0f3F8060FE;
+	fma.rn.ftz.f32 	%f218, %f448, %f449, %f447;
+	fma.rn.ftz.f32 	%f219, %f218, %f449, %f446;
+	fma.rn.ftz.f32 	%f220, %f219, %f449, %f445;
+	fma.rn.ftz.f32 	%f221, %f220, %f449, %f444;
+	fma.rn.ftz.f32 	%f222, %f221, %f449, %f443;
+	fma.rn.ftz.f32 	%f223, %f222, %f449, %f442;
+	neg.f32 	%f224, %f449;
+	selp.f32 	%f225, %f224, %f20, %p12;
+	fma.rn.ftz.f32 	%f450, %f223, %f225, %f225;
+	@%p13 bra 	$L__BB0_16;
+	ex2.approx.ftz.f32 	%f226, %f450;
+	sub.f32 	%f228, %f414, %f226;
+	mov.b32 	%r43, %f228;
+	mov.b32 	%r44, %f20;
+	and.b32  	%r45, %r44, -2147483648;
+	or.b32  	%r46, %r45, %r43;
+	mov.b32 	%f450, %r46;
+$L__BB0_16:
+	.loc	1 0 0
+	cvt.u16.u32 	%rs15, %r17;
+	mov.b32 	%f14, %r23;
+	.loc	1 30 23
+	abs.ftz.f32 	%f77, %f21;
+	setp.ge.f32 	%p15, %f77, 0f3F8060FE;
+	mov.f32 	%f457, 0f3789CA3C;
+	mov.f32 	%f456, 0fB9F560B9;
+	mov.f32 	%f455, 0f3BAC840B;
+	mov.f32 	%f454, 0fBD0C8162;
+	mov.f32 	%f453, 0f3E1CF906;
+	mov.f32 	%f452, 0f3F6A937E;
+	mov.f32 	%f451, 0f3F20D842;
+	mov.f32 	%f458, %f77;
+	@%p15 bra 	$L__BB0_18;
+	mul.f32 	%f458, %f21, %f21;
+	mov.f32 	%f457, 0f38B1E96A;
+	mov.f32 	%f456, 0fBA574D20;
+	mov.f32 	%f455, 0f3BAAD5EA;
+	mov.f32 	%f454, 0fBCDC1BE7;
+	mov.f32 	%f453, 0f3DE718AF;
+	mov.f32 	%f452, 0fBEC093AC;
+	mov.f32 	%f451, 0f3E0375D3;
+$L__BB0_18:
+	.loc	1 0 0
+	cvt.f32.bf16 %r24, %rs15;
+	mul.f32 	%f22, %f14, 0f3F3504F3;
+	.loc	1 30 23
+	setp.ltu.f32 	%p16, %f77, 0f3F8060FE;
+	fma.rn.ftz.f32 	%f243, %f457, %f458, %f456;
+	fma.rn.ftz.f32 	%f244, %f243, %f458, %f455;
+	fma.rn.ftz.f32 	%f245, %f244, %f458, %f454;
+	fma.rn.ftz.f32 	%f246, %f245, %f458, %f453;
+	fma.rn.ftz.f32 	%f247, %f246, %f458, %f452;
+	fma.rn.ftz.f32 	%f248, %f247, %f458, %f451;
+	neg.f32 	%f249, %f458;
+	selp.f32 	%f250, %f249, %f21, %p15;
+	fma.rn.ftz.f32 	%f459, %f248, %f250, %f250;
+	@%p16 bra 	$L__BB0_20;
+	ex2.approx.ftz.f32 	%f251, %f459;
+	sub.f32 	%f253, %f414, %f251;
+	mov.b32 	%r47, %f253;
+	mov.b32 	%r48, %f21;
+	and.b32  	%r49, %r48, -2147483648;
+	or.b32  	%r50, %r49, %r47;
+	mov.b32 	%f459, %r50;
+$L__BB0_20:
+	.loc	1 0 0
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs16}, %r17; }
+	mov.b32 	%f15, %r24;
+	.loc	1 30 23
+	abs.ftz.f32 	%f90, %f22;
+	setp.ge.f32 	%p18, %f90, 0f3F8060FE;
+	mov.f32 	%f466, 0f3789CA3C;
+	mov.f32 	%f465, 0fB9F560B9;
+	mov.f32 	%f464, 0f3BAC840B;
+	mov.f32 	%f463, 0fBD0C8162;
+	mov.f32 	%f462, 0f3E1CF906;
+	mov.f32 	%f461, 0f3F6A937E;
+	mov.f32 	%f460, 0f3F20D842;
+	mov.f32 	%f467, %f90;
+	@%p18 bra 	$L__BB0_22;
+	mul.f32 	%f467, %f22, %f22;
+	mov.f32 	%f466, 0f38B1E96A;
+	mov.f32 	%f465, 0fBA574D20;
+	mov.f32 	%f464, 0f3BAAD5EA;
+	mov.f32 	%f463, 0fBCDC1BE7;
+	mov.f32 	%f462, 0f3DE718AF;
+	mov.f32 	%f461, 0fBEC093AC;
+	mov.f32 	%f460, 0f3E0375D3;
+$L__BB0_22:
+	.loc	1 0 0
+	cvt.f32.bf16 %r25, %rs16;
+	mul.f32 	%f23, %f15, 0f3F3504F3;
+	.loc	1 30 23
+	setp.ltu.f32 	%p19, %f90, 0f3F8060FE;
+	fma.rn.ftz.f32 	%f268, %f466, %f467, %f465;
+	fma.rn.ftz.f32 	%f269, %f268, %f467, %f464;
+	fma.rn.ftz.f32 	%f270, %f269, %f467, %f463;
+	fma.rn.ftz.f32 	%f271, %f270, %f467, %f462;
+	fma.rn.ftz.f32 	%f272, %f271, %f467, %f461;
+	fma.rn.ftz.f32 	%f273, %f272, %f467, %f460;
+	neg.f32 	%f274, %f467;
+	selp.f32 	%f275, %f274, %f22, %p18;
+	fma.rn.ftz.f32 	%f468, %f273, %f275, %f275;
+	@%p19 bra 	$L__BB0_24;
+	ex2.approx.ftz.f32 	%f276, %f468;
+	sub.f32 	%f278, %f414, %f276;
+	mov.b32 	%r51, %f278;
+	mov.b32 	%r52, %f22;
+	and.b32  	%r53, %r52, -2147483648;
+	or.b32  	%r54, %r53, %r51;
+	mov.b32 	%f468, %r54;
+$L__BB0_24:
+	.loc	1 0 0
+	mov.b32 	%f16, %r25;
+	.loc	1 30 23
+	abs.ftz.f32 	%f103, %f23;
+	setp.ge.f32 	%p21, %f103, 0f3F8060FE;
+	mov.f32 	%f475, 0f3789CA3C;
+	mov.f32 	%f474, 0fB9F560B9;
+	mov.f32 	%f473, 0f3BAC840B;
+	mov.f32 	%f472, 0fBD0C8162;
+	mov.f32 	%f471, 0f3E1CF906;
+	mov.f32 	%f470, 0f3F6A937E;
+	mov.f32 	%f469, 0f3F20D842;
+	mov.f32 	%f476, %f103;
+	@%p21 bra 	$L__BB0_26;
+	mul.f32 	%f476, %f23, %f23;
+	mov.f32 	%f475, 0f38B1E96A;
+	mov.f32 	%f474, 0fBA574D20;
+	mov.f32 	%f473, 0f3BAAD5EA;
+	mov.f32 	%f472, 0fBCDC1BE7;
+	mov.f32 	%f471, 0f3DE718AF;
+	mov.f32 	%f470, 0fBEC093AC;
+	mov.f32 	%f469, 0f3E0375D3;
+$L__BB0_26:
+	.loc	1 0 0
+	cvt.u16.u32 	%rs1, %r2;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs2}, %r2; }
+	cvt.u16.u32 	%rs3, %r3;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs4}, %r3; }
+	cvt.u16.u32 	%rs5, %r4;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs6}, %r4; }
+	cvt.u16.u32 	%rs7, %r5;
+	{ .reg .b16 tmp; mov.b32 {tmp, %rs8}, %r5; }
+	mul.f32 	%f24, %f16, 0f3F3504F3;
+	.loc	1 30 23
+	setp.ltu.f32 	%p22, %f103, 0f3F8060FE;
+	fma.rn.ftz.f32 	%f293, %f475, %f476, %f474;
+	fma.rn.ftz.f32 	%f294, %f293, %f476, %f473;
+	fma.rn.ftz.f32 	%f295, %f294, %f476, %f472;
+	fma.rn.ftz.f32 	%f296, %f295, %f476, %f471;
+	fma.rn.ftz.f32 	%f297, %f296, %f476, %f470;
+	fma.rn.ftz.f32 	%f298, %f297, %f476, %f469;
+	neg.f32 	%f299, %f476;
+	selp.f32 	%f300, %f299, %f23, %p21;
+	fma.rn.ftz.f32 	%f477, %f298, %f300, %f300;
+	@%p22 bra 	$L__BB0_28;
+	ex2.approx.ftz.f32 	%f301, %f477;
+	sub.f32 	%f303, %f414, %f301;
+	mov.b32 	%r55, %f303;
+	mov.b32 	%r56, %f23;
+	and.b32  	%r57, %r56, -2147483648;
+	or.b32  	%r58, %r57, %r55;
+	mov.b32 	%f477, %r58;
+$L__BB0_28:
+	.loc	1 0 0
+	cvt.f32.bf16 %r6, %rs1;
+	cvt.f32.bf16 %r7, %rs2;
+	cvt.f32.bf16 %r8, %rs3;
+	cvt.f32.bf16 %r9, %rs4;
+	cvt.f32.bf16 %r10, %rs5;
+	cvt.f32.bf16 %r11, %rs6;
+	cvt.f32.bf16 %r12, %rs7;
+	cvt.f32.bf16 %r13, %rs8;
+	.loc	1 30 23
+	abs.ftz.f32 	%f116, %f24;
+	setp.ge.f32 	%p24, %f116, 0f3F8060FE;
+	mov.f32 	%f484, 0f3789CA3C;
+	mov.f32 	%f483, 0fB9F560B9;
+	mov.f32 	%f482, 0f3BAC840B;
+	mov.f32 	%f481, 0fBD0C8162;
+	mov.f32 	%f480, 0f3E1CF906;
+	mov.f32 	%f479, 0f3F6A937E;
+	mov.f32 	%f478, 0f3F20D842;
+	mov.f32 	%f485, %f116;
+	@%p24 bra 	$L__BB0_30;
+	mul.f32 	%f485, %f24, %f24;
+	mov.f32 	%f484, 0f38B1E96A;
+	mov.f32 	%f483, 0fBA574D20;
+	mov.f32 	%f482, 0f3BAAD5EA;
+	mov.f32 	%f481, 0fBCDC1BE7;
+	mov.f32 	%f480, 0f3DE718AF;
+	mov.f32 	%f479, 0fBEC093AC;
+	mov.f32 	%f478, 0f3E0375D3;
+$L__BB0_30:
+	.loc	1 0 0
+	mov.b32 	%f1, %r6;
+	mov.b32 	%f2, %r7;
+	mov.b32 	%f3, %r8;
+	mov.b32 	%f4, %r9;
+	mov.b32 	%f5, %r10;
+	mov.b32 	%f6, %r11;
+	mov.b32 	%f7, %r12;
+	mov.b32 	%f8, %r13;
+	.loc	1 30 23
+	setp.ltu.f32 	%p25, %f116, 0f3F8060FE;
+	fma.rn.ftz.f32 	%f318, %f484, %f485, %f483;
+	fma.rn.ftz.f32 	%f319, %f318, %f485, %f482;
+	fma.rn.ftz.f32 	%f320, %f319, %f485, %f481;
+	fma.rn.ftz.f32 	%f321, %f320, %f485, %f480;
+	fma.rn.ftz.f32 	%f322, %f321, %f485, %f479;
+	fma.rn.ftz.f32 	%f323, %f322, %f485, %f478;
+	neg.f32 	%f324, %f485;
+	selp.f32 	%f325, %f324, %f24, %p24;
+	fma.rn.ftz.f32 	%f486, %f323, %f325, %f325;
+	@%p25 bra 	$L__BB0_32;
+	ex2.approx.ftz.f32 	%f326, %f486;
+	sub.f32 	%f328, %f414, %f326;
+	mov.b32 	%r59, %f328;
+	mov.b32 	%r60, %f24;
+	and.b32  	%r61, %r60, -2147483648;
+	or.b32  	%r62, %r61, %r59;
+	mov.b32 	%f486, %r62;
+$L__BB0_32:
+	.loc	1 32 18
+	add.f32 	%f345, %f423, 0f3F800000;
+	add.f32 	%f346, %f432, 0f3F800000;
+	add.f32 	%f347, %f441, 0f3F800000;
+	add.f32 	%f348, %f450, 0f3F800000;
+	add.f32 	%f349, %f459, 0f3F800000;
+	add.f32 	%f350, %f468, 0f3F800000;
+	add.f32 	%f351, %f477, 0f3F800000;
+	add.f32 	%f352, %f486, 0f3F800000;
+	.loc	1 35 19
+	mul.f32 	%f353, %f9, %f9;
+	mul.f32 	%f354, %f10, %f10;
+	mul.f32 	%f355, %f11, %f11;
+	mul.f32 	%f356, %f12, %f12;
+	mul.f32 	%f357, %f13, %f13;
+	mul.f32 	%f358, %f14, %f14;
+	mul.f32 	%f359, %f15, %f15;
+	mul.f32 	%f360, %f16, %f16;
+	.loc	1 37 20
+	mul.f32 	%f361, %f353, 0fBF000000;
+	mul.f32 	%f362, %f354, 0fBF000000;
+	mul.f32 	%f363, %f355, 0fBF000000;
+	mul.f32 	%f364, %f356, 0fBF000000;
+	mul.f32 	%f365, %f357, 0fBF000000;
+	mul.f32 	%f366, %f358, 0fBF000000;
+	mul.f32 	%f367, %f359, 0fBF000000;
+	mul.f32 	%f368, %f360, 0fBF000000;
+	.loc	1 38 19
+	mul.f32 	%f330, %f361, 0f3FB8AA3B;
+	ex2.approx.f32 %f329, %f330;
+	mul.f32 	%f332, %f362, 0f3FB8AA3B;
+	ex2.approx.f32 %f331, %f332;
+	mul.f32 	%f334, %f363, 0f3FB8AA3B;
+	ex2.approx.f32 %f333, %f334;
+	mul.f32 	%f336, %f364, 0f3FB8AA3B;
+	ex2.approx.f32 %f335, %f336;
+	mul.f32 	%f338, %f365, 0f3FB8AA3B;
+	ex2.approx.f32 %f337, %f338;
+	mul.f32 	%f340, %f366, 0f3FB8AA3B;
+	ex2.approx.f32 %f339, %f340;
+	mul.f32 	%f342, %f367, 0f3FB8AA3B;
+	ex2.approx.f32 %f341, %f342;
+	mul.f32 	%f344, %f368, 0f3FB8AA3B;
+	ex2.approx.f32 %f343, %f344;
+	.loc	1 40 20
+	mul.f32 	%f369, %f329, 0f3ECC422A;
+	mul.f32 	%f370, %f331, 0f3ECC422A;
+	mul.f32 	%f371, %f333, 0f3ECC422A;
+	mul.f32 	%f372, %f335, 0f3ECC422A;
+	mul.f32 	%f373, %f337, 0f3ECC422A;
+	mul.f32 	%f374, %f339, 0f3ECC422A;
+	mul.f32 	%f375, %f341, 0f3ECC422A;
+	mul.f32 	%f376, %f343, 0f3ECC422A;
+	.loc	1 41 19
+	mul.f32 	%f377, %f9, %f369;
+	mul.f32 	%f378, %f10, %f370;
+	mul.f32 	%f379, %f11, %f371;
+	mul.f32 	%f380, %f12, %f372;
+	mul.f32 	%f381, %f13, %f373;
+	mul.f32 	%f382, %f14, %f374;
+	mul.f32 	%f383, %f15, %f375;
+	mul.f32 	%f384, %f16, %f376;
+	.loc	1 42 20
+	fma.rn.f32 	%f385, %f345, 0f3F000000, %f377;
+	fma.rn.f32 	%f386, %f346, 0f3F000000, %f378;
+	fma.rn.f32 	%f387, %f347, 0f3F000000, %f379;
+	fma.rn.f32 	%f388, %f348, 0f3F000000, %f380;
+	fma.rn.f32 	%f389, %f349, 0f3F000000, %f381;
+	fma.rn.f32 	%f390, %f350, 0f3F000000, %f382;
+	fma.rn.f32 	%f391, %f351, 0f3F000000, %f383;
+	fma.rn.f32 	%f392, %f352, 0f3F000000, %f384;
+	.loc	1 43 19
+	mul.f32 	%f393, %f1, %f385;
+	mul.f32 	%f394, %f2, %f386;
+	mul.f32 	%f395, %f3, %f387;
+	mul.f32 	%f396, %f4, %f388;
+	mul.f32 	%f397, %f5, %f389;
+	mul.f32 	%f398, %f6, %f390;
+	mul.f32 	%f399, %f7, %f391;
+	mul.f32 	%f400, %f8, %f392;
+	.loc	1 45 40
+	mov.b32 	%r63, %f393;
+	cvt.rn.bf16.f32 %rs17, %r63;
+	mov.b32 	%r64, %f394;
+	cvt.rn.bf16.f32 %rs18, %r64;
+	mov.b32 	%r65, %f395;
+	cvt.rn.bf16.f32 %rs19, %r65;
+	mov.b32 	%r66, %f396;
+	cvt.rn.bf16.f32 %rs20, %r66;
+	mov.b32 	%r67, %f397;
+	cvt.rn.bf16.f32 %rs21, %r67;
+	mov.b32 	%r68, %f398;
+	cvt.rn.bf16.f32 %rs22, %r68;
+	mov.b32 	%r69, %f399;
+	cvt.rn.bf16.f32 %rs23, %r69;
+	mov.b32 	%r70, %f400;
+	cvt.rn.bf16.f32 %rs24, %r70;
+	mov.b32 	%r75, {%rs17, %rs18};
+	mov.b32 	%r76, {%rs19, %rs20};
+	mov.b32 	%r77, {%rs21, %rs22};
+	mov.b32 	%r78, {%rs23, %rs24};
+	@%p1 st.global.v4.b32 [ %rd7 + 0 ], { %r75, %r76, %r77, %r78 };
+	.loc	1 45 4
+	ret;
+$L__tmp1:
+$L__func_end0:
+
+}
+	// .globl	__nv_erff
+.visible .func  (.param .b32 func_retval0) __nv_erff(
+	.param .b32 __nv_erff_param_0
+)
+{
+	.reg .pred 	%p<4>;
+	.reg .b32 	%r<5>;
+	.reg .f32 	%f<49>;
+$L__func_begin1:
+
+	ld.param.f32 	%f14, [__nv_erff_param_0];
+	abs.ftz.f32 	%f1, %f14;
+	setp.ge.f32 	%p1, %f1, 0f3F8060FE;
+	mov.f32 	%f46, 0f3789CA3C;
+	mov.f32 	%f45, 0fB9F560B9;
+	mov.f32 	%f44, 0f3BAC840B;
+	mov.f32 	%f43, 0fBD0C8162;
+	mov.f32 	%f42, 0f3E1CF906;
+	mov.f32 	%f41, 0f3F6A937E;
+	mov.f32 	%f40, 0f3F20D842;
+	mov.f32 	%f47, %f1;
+	@%p1 bra 	$L__BB1_2;
+	mul.f32 	%f47, %f14, %f14;
+	mov.f32 	%f46, 0f38B1E96A;
+	mov.f32 	%f45, 0fBA574D20;
+	mov.f32 	%f44, 0f3BAAD5EA;
+	mov.f32 	%f43, 0fBCDC1BE7;
+	mov.f32 	%f42, 0f3DE718AF;
+	mov.f32 	%f41, 0fBEC093AC;
+	mov.f32 	%f40, 0f3E0375D3;
+$L__BB1_2:
+	setp.ltu.f32 	%p2, %f1, 0f3F8060FE;
+	fma.rn.ftz.f32 	%f29, %f46, %f47, %f45;
+	fma.rn.ftz.f32 	%f30, %f29, %f47, %f44;
+	fma.rn.ftz.f32 	%f31, %f30, %f47, %f43;
+	fma.rn.ftz.f32 	%f32, %f31, %f47, %f42;
+	fma.rn.ftz.f32 	%f33, %f32, %f47, %f41;
+	fma.rn.ftz.f32 	%f34, %f33, %f47, %f40;
+	neg.f32 	%f35, %f47;
+	selp.f32 	%f36, %f35, %f14, %p1;
+	fma.rn.ftz.f32 	%f48, %f34, %f36, %f36;
+	@%p2 bra 	$L__BB1_4;
+	ex2.approx.ftz.f32 	%f37, %f48;
+	mov.f32 	%f38, 0f3F800000;
+	sub.f32 	%f39, %f38, %f37;
+	mov.b32 	%r1, %f39;
+	mov.b32 	%r2, %f14;
+	and.b32  	%r3, %r2, -2147483648;
+	or.b32  	%r4, %r3, %r1;
+	mov.b32 	%f48, %r4;
+$L__BB1_4:
+	st.param.f32 	[func_retval0+0], %f48;
+	ret;
+$L__func_end1:
+
+}
+	.file	1 "/tmp/torchinductor_root/5j/c5jxaguxho3nhrlt5vcinnz5fevodumlpwn4wyb2vx3xrveicerl.py"
+	.section	.debug_abbrev
+	{
+.b8 1
+.b8 17
+.b8 1
+.b8 37
+.b8 8
+.b8 19
+.b8 5
+.b8 3
+.b8 8
+.b8 16
+.b8 6
+.b8 27
+.b8 8
+.b8 180
+.b8 66
+.b8 12
+.b8 17
+.b8 1
+.b8 18
+.b8 1
+.b8 0
+.b8 0
+.b8 2
+.b8 46
+.b8 0
+.b8 17
+.b8 1
+.b8 18
+.b8 1
+.b8 64
+.b8 10
+.b8 135
+.b8 64
+.b8 8
+.b8 3
+.b8 8
+.b8 58
+.b8 11
+.b8 59
+.b8 11
+.b8 63
+.b8 12
+.b8 0
+.b8 0
+.b8 0
+	}
+	.section	.debug_info
+	{
+.b32 176
+.b8 2
+.b8 0
+.b32 .debug_abbrev
+.b8 8
+.b8 1
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 0
+.b8 2
+.b8 0
+.b8 99
+.b8 53
+.b8 106
+.b8 120
+.b8 97
+.b8 103
+.b8 117
+.b8 120
+.b8 104
+.b8 111
+.b8 51
+.b8 110
+.b8 104
+.b8 114
+.b8 108
+.b8 116
+.b8 53
+.b8 118
+.b8 99
+.b8 105
+.b8 110
+.b8 110
+.b8 122
+.b8 53
+.b8 102
+.b8 101
+.b8 118
+.b8 111
+.b8 100
+.b8 117
+.b8 109
+.b8 108
+.b8 112
+.b8 119
+.b8 110
+.b8 52
+.b8 119
+.b8 121
+.b8 98
+.b8 50
+.b8 118
+.b8 120
+.b8 51
+.b8 120
+.b8 114
+.b8 118
+.b8 101
+.b8 105
+.b8 99
+.b8 101
+.b8 114
+.b8 108
+.b8 46
+.b8 112
+.b8 121
+.b8 0
+.b32 .debug_line
+.b8 47
+.b8 116
+.b8 109
+.b8 112
+.b8 47
+.b8 116
+.b8 111
+.b8 114
+.b8 99
+.b8 104
+.b8 105
+.b8 110
+.b8 100
+.b8 117
+.b8 99
+.b8 116
+.b8 111
+.b8 114
+.b8 95
+.b8 114
+.b8 111
+.b8 111
+.b8 116
+.b8 47
+.b8 53
+.b8 106
+.b8 0
+.b8 1
+.b64 $L__func_begin0
+.b64 $L__func_end0
+.b8 2
+.b64 $L__func_begin0
+.b64 $L__func_end0
+.b8 1
+.b8 156
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 95
+.b8 95
+.b8 48
+.b8 100
+.b8 49
+.b8 100
+.b8 50
+.b8 100
+.b8 101
+.b8 0
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 95
+.b8 95
+.b8 48
+.b8 100
+.b8 49
+.b8 100
+.b8 50
+.b8 100
+.b8 101
+.b8 0
+.b8 1
+.b8 18
+.b8 1
+.b8 0
+	}
+	.section	.debug_pubnames
+	{
+.b32 $L__pubNames_end0-$L__pubNames_start0
+$L__pubNames_start0:
+.b8 2
+.b8 0
+.b32 .debug_info
+.b32 180
+.b32 125
+.b8 116
+.b8 114
+.b8 105
+.b8 116
+.b8 111
+.b8 110
+.b8 95
+.b8 95
+.b8 48
+.b8 100
+.b8 49
+.b8 100
+.b8 50
+.b8 100
+.b8 101
+.b8 0
+.b32 0
+$L__pubNames_end0:
+	}
+	.section	.debug_pubtypes
+	{
+.b32 $L__pubTypes_end0-$L__pubTypes_start0
+$L__pubTypes_start0:
+.b8 2
+.b8 0
+.b32 .debug_info
+.b32 180
+.b32 0
+$L__pubTypes_end0:
+	}
+	.section	.debug_loc	{	}

.triton/dump/a4652f539404a11e3c068d96115a7427/triton_.ttir

@@ -0,0 +1,18 @@
+module {
+  tt.func public @triton__0d1d2de(%arg0: !tt.ptr<bf16, 1> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<f32, 1> {tt.divisibility = 16 : i32}, %arg2: i32 {tt.divisibility = 16 : i32, tt.max_divisibility = 16 : i32}) attributes {noinline = false} {
+    %c256_i32 = arith.constant 256 : i32
+    %0 = tt.get_program_id x : i32
+    %1 = arith.muli %0, %c256_i32 : i32
+    %2 = tt.make_range {end = 256 : i32, start = 0 : i32} : tensor<256xi32>
+    %3 = tt.splat %1 : (i32) -> tensor<256xi32>
+    %4 = arith.addi %3, %2 : tensor<256xi32>
+    %5 = tt.splat %arg0 : (!tt.ptr<bf16, 1>) -> tensor<256x!tt.ptr<bf16, 1>>
+    %6 = tt.addptr %5, %4 : tensor<256x!tt.ptr<bf16, 1>>, tensor<256xi32>
+    %7 = tt.load %6 {cache = 1 : i32, evict = 1 : i32, isVolatile = false} : tensor<256xbf16>
+    %8 = arith.extf %7 : tensor<256xbf16> to tensor<256xf32>
+    %9 = tt.splat %arg1 : (!tt.ptr<f32, 1>) -> tensor<256x!tt.ptr<f32, 1>>
+    %10 = tt.addptr %9, %4 : tensor<256x!tt.ptr<f32, 1>>, tensor<256xi32>
+    tt.store %10, %8 {cache = 1 : i32, evict = 1 : i32} : tensor<256xf32>
+    tt.return
+  }
+}