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.

Sign up for GitHub

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

Jump to bottom

Minor Additions to Enable Tiling and Explicit Memory Movement Transformations #1636

Merged
merged 10 commits into from
Nov 29, 2024
Merged

Minor Additions to Enable Tiling and Explicit Memory Movement Transformations #1636

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
7149621
Add options to force a syncthread within a map and prefer the gpu_blo…
ThrudPrimrose Sep 3, 2024
8efe6b5
Add skew option map tiling transformation
ThrudPrimrose Sep 3, 2024
e8c134a
Improve grammar
ThrudPrimrose Sep 3, 2024
675e1dc
Incorporate step size to thread block computation
ThrudPrimrose Sep 5, 2024
dc6fea3
Merge branch 'master' into additions
ThrudPrimrose Sep 5, 2024
5e21392
The step is not really necessary for the block size detection if the …
ThrudPrimrose Sep 5, 2024
cd61594
Fix typo
ThrudPrimrose Sep 5, 2024
40de8f6
Add gpu force syncthreads to to generate MapExit
ThrudPrimrose Sep 5, 2024
a7ab6c3
Merge branch 'spcl:master' into additions
ThrudPrimrose Oct 15, 2024
cac7515
Merge branch 'main' into additions
ThrudPrimrose Nov 21, 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
3 changes: 3 additions & 0 deletions dace/codegen/targets/cuda.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2620,6 +2620,9 @@ def _generate_NestedSDFG(self, sdfg: SDFG, cfg: ControlFlowRegion, dfg: StateSub

def _generate_MapExit(self, sdfg: SDFG, cfg: ControlFlowRegion, dfg: StateSubgraphView, state_id: int,
node: nodes.MapExit, function_stream: CodeIOStream, callsite_stream: CodeIOStream) -> None:
if isinstance(node, nodes.MapExit) and node.map.gpu_force_syncthreads:
callsite_stream.write('__syncthreads();', cfg, state_id)

if node.map.schedule == dtypes.ScheduleType.GPU_Device:
# Remove grid invocation conditions
for i in range(len(node.map.params)):
Expand Down
2 changes: 2 additions & 0 deletions dace/sdfg/nodes.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -930,6 +930,8 @@ class Map(object):
"(including tuples) sets it explicitly.",
serialize_if=lambda m: m.schedule in dtypes.GPU_SCHEDULES)

gpu_force_syncthreads = Property(dtype=bool, desc="Force a call to the __syncthreads for the map", default=False)

def __init__(self,
label,
params,
Expand Down
4 changes: 4 additions & 0 deletions dace/transformation/dataflow/tiling.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,6 +33,8 @@ class MapTiling(transformation.SingleStateTransformation):
divides_evenly = Property(dtype=bool, default=False, desc="Tile size divides dimension length evenly")
tile_trivial = Property(dtype=bool, default=False, desc="Tiles even if tile_size is 1")

skew = Property(dtype=bool, default=False, desc="If True, offsets inner tile back such that it starts with zero")

@staticmethod
def annotates_memlets():
return True
Expand Down Expand Up @@ -92,6 +94,7 @@ def apply(self, graph: SDFGState, sdfg: SDFG):
stripmine.tile_stride = str(tile_stride)
stripmine.divides_evenly = True
stripmine.tile_offset = str(offset)
stripmine.skew = self.skew
stripmine.apply(graph, sdfg)
removed_maps += 1
else:
Expand All @@ -101,6 +104,7 @@ def apply(self, graph: SDFGState, sdfg: SDFG):
stripmine.tile_stride = str(tile_stride)
stripmine.divides_evenly = self.divides_evenly
stripmine.tile_offset = str(offset)
stripmine.skew = self.skew
stripmine.apply(graph, sdfg)

# apply to the new map the schedule of the original one
Expand Down