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

Refactor: Prerequisite for AutoScheduler #362

Merged
merged 57 commits into from
Dec 22, 2024
Merged
Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
57 commits
Select commit Hold shift + click to select a range
6782d42
gemm
hikettei Dec 20, 2024
44a490f
Tile
hikettei Dec 20, 2024
610316a
search tile
hikettei Dec 20, 2024
0cd518a
Merge branch 'main' into opt-gemm
hikettei Dec 20, 2024
f6bcc57
OpenBLAS
hikettei Dec 20, 2024
253cf6f
90GFLops
hikettei Dec 20, 2024
d4c08c4
wip
hikettei Dec 20, 2024
db52790
wip
hikettei Dec 20, 2024
082ae9d
wip
hikettei Dec 20, 2024
9446987
tile param search
hikettei Dec 20, 2024
6a92698
finish an experiment
hikettei Dec 20, 2024
0e435c1
apply tiling without causing segv
hikettei Dec 20, 2024
e32cec2
tiling all bands
hikettei Dec 20, 2024
c4215bc
things are moved
hikettei Dec 20, 2024
4c2de70
move everything into auto-scheduler
hikettei Dec 21, 2024
785d2eb
asd
hikettei Dec 21, 2024
34656d2
fix: tiling-size related gc error
hikettei Dec 21, 2024
6a0c211
An initial attempt to unroll: tile based
hikettei Dec 21, 2024
0ab186d
moved
hikettei Dec 21, 2024
2ce3075
split
hikettei Dec 21, 2024
b22d20f
.
hikettei Dec 21, 2024
874cd17
moved isl things to auto-scheduler
hikettei Dec 21, 2024
b47fd09
refactor
hikettei Dec 21, 2024
e92725d
synchronize baseline
hikettei Dec 21, 2024
e1e5dc7
updt
hikettei Dec 21, 2024
dfca38d
updt
hikettei Dec 21, 2024
0ed2740
typo
hikettei Dec 21, 2024
0adb742
typo
hikettei Dec 21, 2024
378c310
typo
hikettei Dec 21, 2024
d2de5a1
Tweak
hikettei Dec 21, 2024
cd01fe9
nested unroll directive
hikettei Dec 21, 2024
64b8c7c
Unroll for static and simple case
hikettei Dec 21, 2024
3ac9213
Fold MOD
hikettei Dec 21, 2024
0d4095f
nope
hikettei Dec 21, 2024
642acc9
Initial attempt of unrolling
hikettei Dec 21, 2024
ea829fd
copy and remove UNROLL_PARENT for the reminder
hikettei Dec 21, 2024
b7ba19f
Loop Unrolling is worked
hikettei Dec 21, 2024
6df4461
Unrolling softmax
hikettei Dec 21, 2024
994a615
wip: generating a sketch
hikettei Dec 21, 2024
ca032f8
Identify the sketch
hikettei Dec 21, 2024
4cdec95
TODO
hikettei Dec 21, 2024
1c5096e
remove extra unroll reminder
hikettei Dec 21, 2024
bfb2d11
coincident for elwise
hikettei Dec 22, 2024
771fa79
fix: unroll
hikettei Dec 22, 2024
aad075a
quickload
hikettei Dec 22, 2024
1b56342
MemoryPlanner should produce the same result
hikettei Dec 22, 2024
e77703a
Fix for unroll reminder
hikettei Dec 22, 2024
0e9fff8
try unrolling tiled bands
hikettei Dec 22, 2024
6674d7b
ignore tiled band dims
hikettei Dec 22, 2024
ce60a52
base-item
hikettei Dec 22, 2024
5ba4fa3
remove: caten/polyhedral
hikettei Dec 22, 2024
26a4a66
clean up
hikettei Dec 22, 2024
e63ac07
rem dep
hikettei Dec 22, 2024
546d0c4
caten instead
hikettei Dec 22, 2024
f73c49f
updt
hikettei Dec 22, 2024
04cb131
updt
hikettei Dec 22, 2024
c832f5e
updt
hikettei Dec 22, 2024
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
Next Next commit
try unrolling tiled bands
  • Loading branch information
hikettei committed Dec 22, 2024
commit 0e9fff8fd80c603e3180f93a986cea91f8960cd4
3 changes: 2 additions & 1 deletion source/codegen/auto-scheduler/ast-parser.lisp
Original file line number Diff line number Diff line change
Expand Up @@ -73,6 +73,7 @@ scop.lisp for the opposite things.
;; UNROLL_BODY is triggered by the UNROLL_PARENT. Without it the form is ignored.
(assert (null (astfor-marks user)) () "UNROLL_BODY should be orthogonal with other directives.")
(setf (astfor-marks user) (list mark)))
((equalp mark "TILE_BAND"))
(T
(warn "mark: ignored the mark ~a for ~a" mark user))))
(otherwise
Expand Down Expand Up @@ -207,7 +208,7 @@ scop.lisp for the opposite things.
((ASTBlock :body body) (map 'list #'lower body))
((AstFor :idx idx :from upfrom :to to :by by :body body :scope scope)
;; remove an empty loop
(when (not (expr-scalar-equivalent-p upfrom (expr-detach-loop-bound to)))
g (when (not (expr-scalar-equivalent-p upfrom (expr-detach-loop-bound to)))
(push idx space)
(push (r/for idx upfrom to by scope) new-graph)
(lower body)
Expand Down
3 changes: 2 additions & 1 deletion source/codegen/auto-scheduler/auto-scheduler.lisp
Original file line number Diff line number Diff line change
Expand Up @@ -69,10 +69,11 @@ So we are going to apply an optiimzation method which takes a long time to searc
(print "REUSE")
;; Data Reuse kernel has a chance to apply the tiling
;; And Tuning the tiling size tested by the measurer
;; (caten/codegen/tiling::apply-tile (getattr node :polyhedral) 16)
(caten/codegen/tiling::apply-tile (getattr node :polyhedral) 64) ;; Tile_Size = Optimal_Tile_Size * N_UNROLL
(caten/codegen/unroll:apply-unroll node 4)
;; (caten/codegen/unroll::apply-unroll node 4)
;; [TODO] After Tiling, Unroll the outermost tile band.
;; [TODO] If the outermost loop is a tile -> #pragma omp parallel for collapse(2)
(caten/codegen/coincidence:apply-parallel
(getattr node :polyhedral)
(auto-scheduler-n-global-loops auto-scheduler)))
Expand Down
6 changes: 4 additions & 2 deletions source/codegen/auto-scheduler/tiling.lisp
Original file line number Diff line number Diff line change
Expand Up @@ -60,8 +60,10 @@ dims is used to specify the tiling sizes for each dimension."
(shift-band-zero band)
(let* ((tiling-sizes (tiling-sizes band :size-default size-default :dims dims))
(partial-schedule (tile-partial-schedule partial-schedule tiling-sizes))
(tiled-sched (multi-union-pw-aff-add partial-schedule shift)))
(schedule-node-insert-partial-schedule band tiled-sched))))
(tiled-sched (multi-union-pw-aff-add partial-schedule shift))
(tiled-sched (schedule-node-insert-partial-schedule band tiled-sched))
(tiled-sched (schedule-node-insert-mark tiled-sched (isl::make-id-from-str "TILE_BAND"))))
tiled-sched)))

(defun get-tileable-bands (poly)
(map-schedule-nodes
Expand Down
4 changes: 3 additions & 1 deletion source/codegen/auto-scheduler/unroll.lisp
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,9 @@

(defun mark-unroll-p (mark)
(declare (type string mark))
(scan "UNROLL" (string-upcase mark)))
(or
(scan "TILE_BAND" (string-upcase mark))
(scan "UNROLL" (string-upcase mark))))

(defun mark-unroll-body-p (mark) (scan "UNROLL_BODY" (string-upcase mark)))
(defun mark-unroll-parent-p (mark) (scan "UNROLL_PARENT" (string-upcase mark)))
Expand Down