Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Generated Backend #425

Closed
wants to merge 4 commits into from
Closed

Generated Backend #425

wants to merge 4 commits into from

Conversation

micmelesse
Copy link
Collaborator

@micmelesse micmelesse commented Dec 13, 2023

This pr does the following

  • updates the old backend to a new generated backend
    • this will make maintance much easier
  • adds a tool (scripts/amd/gen_hip_backend.py) used to generate the backend
    • the tool is based on the parsing lib tree-sitter
    • it parses the c++, python, cmake and tablgen files in this fork and modifies them to be used for the backend
    • the modifications involve altering Macros (defines, includes), imports, types, functions, namespaces, strings.
  • adds a github actions tests for both upstream testing and fork testing
    • in the future this will allows us to move away from jenkins
    • fork test
      • is the same as the jenkins ci
    • upstream test
      • clones the upstream
      • checkouts it out to the last shared commit with upstream
      • applies a patch
      • checks out this branch as a backend
      • builds
      • runs unit tests
  • when reviewing you can ignore python/third_party/hip/gen. It is the generated files from the rest of the fork. If the upstream test works It should be fine.

@scxiao
Copy link

scxiao commented Dec 13, 2023

https://github.com/ROCmSoftwarePlatform/triton/blob/3538df7d4cddbc0770c52bffef17e25ad68792e4/python/triton/third_party/hip/hip_backend.py#L580
It cannot hard coded to 64 since Navi31 has wave size 32. The PR for dynamic warp size is at:

https://github.com/ROCmSoftwarePlatform/triton/pull/244

@micmelesse micmelesse force-pushed the gen_backend_minimal branch 6 times, most recently from da42057 to 3f09070 Compare December 21, 2023 17:44
Gen Backend with CI for Fork and Upstream Tests

This is a combination of 58 commits.

Gen Backend with CI

Gen Backend that works with GEN folder

Gen Backend that works with GEN folder

This is a combination of 8 commits.

Gen Backend that works

This is a combination of 3 commits.

Gen Backend that works

This is a combination of 170 commits.

Buildable backend

This is a combination of 73 commits.

Gen HipBackend

This is a combination of 48 commits.

gen TritonGPUROCMDialect

This is a combination of 2 commits.

Gen TritonGPUROCMDialect

This is a combination of 33 commits.

add scripts

remove hip dir

check and remove hip folder

save ref code

chmod 777 copied files

traverse asts for files

process everything

print preproc directives

edit def and ifdef but second edit corrupt

good edits

edit defines

add clang parse

add lib ref

rewrite includes

add rewrite namespace function

reparse code after rewrite

rewrite use_rocm

rework analysis folder

fix types.cpp error

add debug flag

fix bug in rewrite_namespace. seperate out types rewrite

just replace gpu:: to gpu_rocm for now

add cmake parsing

select file to process

add cmake rewrites

add tablegen

translate TritonGPUDialect

rewrite tablen include, class and parent classes

change #ifdefs in tablegen code strings

map functions and types in cpp

split up tablegen rewrites

search tablegen files for MAP entries

split translation map by language and type of translation

clean

rewrite cpp define follows same template

new TritonGPUROCM IR

This is a combination of 3 commits.

add main conversion pass and dialect transform

remove conversions

new TritonGPUROCM IR

rewrite specific file

fix dir rewrites

use AMD instead of ROCM

fix cmake bug

back to GPUROCM for now

throw exceptsion on failed rewrite

rewrite namespace_alias_definition

fixes

improve tablegen string rewrite

new IR/Transforms headers

Debug Lists, Flags, Fix Namespace Issue, Pretty Print

fix utility.h bug

good transform files

can use paths from root now

add lang to cpp rewrite funcs

rewrite lambdas

black and add_edit function

simplify namespace renames

rewrite error nodes

more transforms

transforms done

good cmake for transform

add Conversion Ref

good conversion cmake

TritonToTrtionGPU stuff 1

tritontotpgu header

triton to tritongpu conversion

some fixes

clean up namespace

good ttir to ttgir cpp

ttgir conversion pass done

bring TritonGPUROCMToLLVM

TritonGPUToLLVM headers working

work TritonGPUToLLVM cpp files

save good cpp

add second to last cpp

cpp namespace function runs at the end

fix double rewrites

finish conversion passes

add AnalysisROCM from triton-mlir

good Analysis

good analyisi

rewrite field declarations

save

debug target runs on debug

update includes

save before refactor

refactor apply edits stage

restructre code remove duplication

fix utiltly.h

put in everything back

remove chmod and minor issues

remove everything

bring back triton-mlir

gen_backend

remove translate map stage in cpp. USe in specific rewrite functions

format script

headers are good

remove double rewrites in HSACO

run all

make script use gitpython and also run relative to current dir

save cmake

chmod rewritten file

strict check on for cmake args

chmod python/triton/third_party/hip/CMakeLists.txt

define TRITONGPU_DEFAULT_WARPSIZE

chmod works on file

add argparse

create src -> dst map

smart dst generation works

gen abs paths

fix issue with path rewrite

fix git clone issue

map ASMBuilder and MLIRGPUops cmake modules

rewrite mlir::triton::gpu also

just make ifu level upstream commit

ignore TMAMetadataTy

rewrite base class

add path arg

replace TritonGPU in using declrations

narrow TMAMetadataTy ignore

fix rewrite issues with Fields

get original source path

throw error if path gen fails

rewrite return types

rewrite function params

general class rewrite

check every rewrite

more subtle rewrite check that use rewrite tuple

simplify namespace rewrite

smart namespace rewrite

simplify type rewrites

create traverse_node_and_exec_fn

better valid namespace check

rewrite functions simplified

rewrite namespaces in fields

fix bug with function args rewrite

rewrite namespace defintions

rewrite calls in functions

fix bug with nested namespaces

nice

unify range checks

traverse only qualified_identifier in namespace rewrite

rewrite namespace in function_declarators

deal with namespace_alias_definition

rewrite qualified function names

namespace rewrite in macro args

use triton:: for any tritongpu usage

simpler is_in_range check and add special rule for bad file Allocation.cpp

field_initializer can start with template method

remove print

ignore TMAInfo

rename addExternalLibs to addExternalLibsROCM

rewrite the args

simplify rewrite_cmake_arguments

New Name for Docs

try again

pass rewrite map as arg

make sure source is utf-8

change triton.cc encoding

add filter function

rename

check if it builds

declar function

use rewrite map from json, gen rocm_backend_for_triton.cc

merge language and file rewrite rules

run everything

format script

save triton_rocm.cc

fix path bugs

disable llvmir import

nice print

save

remove isrocm with out removing comma

remove comma

ignore createTritonGPURewriteTensorPointerPass

add archinfo

rewrite define

gen SysROCM

copy hsa headers

remove print

gen LLVMROCMIR

improve LLVMROCMIR gen and also unify defines

fix include bug

add include for l

comment out createConvertNVGPUToLLVMPass

add TritonLLVMROCMIR in build step

comment out code that is not needed

try renaming

update

comment out init_triton_runtime

comment out stuff

fix import

create hip_bindings.py

gen hip_bindings.py

save compiler.py change

rewrite python bindings

change import

fix relative import

save

remove bad imports

remove ref to runtime

remove translate_llvmir_to_ptx

rewrite relative imports

copy gpu_matrix_core_version

get new context

remove arith dialect

remove everything except TritonGPUROCMDialect

add LLVM dialect

update

comment out registeration

comment out regisreation

empty context

gen TritonROCM

rewrite headers

use TritonGPU

fix build bug

use TTROCM_

fix tablegen

fix tablegen include bug

fix more bugs

handle path bug

rewrite tablegen values

fix no ascii chars

fix allocation.cpp asci issue

fix TritonOps ascii char bug

fix TritonDialect bug

tablegen deps are TritonROCMDialect

fix asci bug in TritonGPUops

use triton_rocm namespace

more triton_rocm

rewrite tablegen dialect strictly

seperate out comment out

gen PointerType correctly

change cpp namespaces

remove triton:: in td files

skip nvgpu

save mp code

rename trait functions

fix in tablegen

try again

rewrite friends

add MLIRTypes

use MLIRTypesROCM

rewrite preproc fund def

more bug fixes

rewrite in for range

nvgpu enable again

back

rewrite identifer in qualifed id

add base class fix

fix tablegen string bug

user triton for nv_gpu

rewrite on pointer exps

fix bug

add TritonNvidiaGPU dialect

start tritonnv stuff

add nvgpu

just gen everything

gen NVGPUROCM

fix TritonNvidiaGPUROCMTransforms

include

save

bug fix

save

save

fix more

rewrite TritonNVGPU funcs

tqdm

create gen

expand more qua_identifer rewrites

rewrite preproc args

fix more qi issue

fix cmake module bugs

rewrite_map is all files

save cpp progress

add filter arg

working td

good td

better td

clean up tablgen string rewrites

add PTX aswell!

gen ptxrocm

gen all

fix bug

fix symbol issue

just pass everything

gen code_generator_rocm

gen more python

try

gen python folders

rewrite dirs with dest in map

test it out

add more files

uncomment out all functions

remove inits

add everything

add inits back

commend out get_backend imports

rewrite imports

add print

add dialects

move add_plan_cta_pass to cuda

gen debug output

create new context!

remove print

rename types and imports for py

get src path from dest path

better import rewrites

parallel flag

minor bugs

working py imports?

rename py strings loss

chmod 777 everything

fix namespace reg issue. Just ttir to file and then parse again

comment out prints

triton_gpu_rocm. to triton_gpu.

remove ttir stuff

use temp files

move shift code to backend

add shift ops to the right file

update gened code

min diff

update gen

convert clusterinfo

get attrs

update gened files

just use _triton

remove gened stuff

save

use gen folder

rewrite filenames

rewrite cmake files correctly

fix relative import

fix more import errors

regen backend

regen after IFU

rm -rf tree_sitter repo

gen include, lib, python

chmod 660

fix namespace bug

fix python issues

fix bug

skip test_bin_op_constexpr tests

fix bug

remove prints

remove prints

clean up scripts

clean up HSACOTranslation.h

update gen

add upstream test

fix bug

cancel in progress

run on rocm

restore offline tests

skip pre_commit for now and add precommit script

add script to diff upstream

update tests

fix bugs

just one test

minor bugs

deal with empty path

minor bug

clone submodules on upstream test

bug with diff

Check out Shared Commit

minor bug

use my fork of upstream

minor bug

keep name fo diffs

try again

show commit and dir

pwd

no upstream test

install torch to detect ROCM

run everything

query warp size

have default for queried warp size

skip test_gemm failure

skip test_print bug

match jenkins ci

just test_core_amd.py

clean everything before test

don't clear /tmp

print archinfo

Update hip_backend to use libhsa-runtime for arch info,

brings in path changes for pytorch triton wheels

print failure

remove try block

print getArchInfo call

minor bug

fail if arch_info is none

docker action

disable fork test

run on pr

try hi

leave it for now

gracefully checkout

fix minor bug

wrapper for get_arch_info

run all of test_core_amd

add test_subprocess

add build upstream script

update gen

update build_upstream script

check for TRITON_CODEGEN_AMD_HIP_BACKEND

clean up build upstream script

update upstream build step

fix cmake if bug

print build upstream step

skip test_cas

fetch before backend checkout

just run upstream test

fetch all

checkout branch name

use github env variable

try head ref

add set branch name

show git log

use patch

use absolute path for script path

git status

try with quotes

install on system

run and test

pick python

remove workerss

try personal branch

fix bug

don't checkout

just run pytest

try what they do

test fork

update backend

add check step

bug!

check after build

use root dir

use venv

activate env before each step

install wheel pacakge

add pytest

more deps

use venv for fork tests

try again

try just pytest

install pytest-xdist

try new patch

new patch
@micmelesse micmelesse closed this Jan 11, 2024
@micmelesse
Copy link
Collaborator Author

moveing the work on the ci to #455. Generating the backend might not be needed.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

2 participants