Skip to content

Commit

Permalink
Merge branch ylee/94-problems-using-milhoja-with-a-new-taskfunction i…
Browse files Browse the repository at this point in the history
…nto kweide/94+96-current
  • Loading branch information
kweide committed Aug 26, 2024
2 parents 17ca37f + 3ece389 commit 6c1e22a
Show file tree
Hide file tree
Showing 13 changed files with 104 additions and 133 deletions.
28 changes: 0 additions & 28 deletions includes/Milhoja_IntVect.h
Original file line number Diff line number Diff line change
Expand Up @@ -109,20 +109,6 @@ class IntVect
#endif
}

//! Return second element of vector, or base if MILHOJA_NDIM<2
#ifdef ACC_ROUTINE_FOR_METH
#pragma acc routine seq
#endif
int JJ(const int base) const {
#if (MILHOJA_NDIM == 1)
return base;
#elif (MILHOJA_NDIM == 2) || (MILHOJA_NDIM == 3)
return j_;
#else
#error "MILHOJA_NDIM not defined or invalid"
#endif
}

//! Return third element of vector, or 0 if MILHOJA_NDIM<3
#ifdef ACC_ROUTINE_FOR_METH
#pragma acc routine seq
Expand All @@ -137,20 +123,6 @@ class IntVect
#endif
}

//! Return third element of vector, or base if MILHOJA_NDIM<3
#ifdef ACC_ROUTINE_FOR_METH
#pragma acc routine seq
#endif
int KK(const int base) const {
#if (MILHOJA_NDIM == 1) || (MILHOJA_NDIM == 2)
return base;
#elif (MILHOJA_NDIM == 3)
return k_;
#else
#error "MILHOJA_NDIM not defined or invalid"
#endif
}

//! Get and set values of the internal array.
int& operator[] (const unsigned int i) {
switch (i) {
Expand Down
10 changes: 5 additions & 5 deletions src/Milhoja_TileFlashxr.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -34,13 +34,13 @@ TileFlashxr::TileFlashxr(const FlashxrTileRawPtrs tP,
const FlashxTileRawInts tI,
const FlashxTileRawReals tR)
: Tile{},
level_{static_cast<unsigned int>(tI.level - 1)}, // 1-based Flash-X convention -> 0-based
level_{static_cast<unsigned int>(tI.level - 1)}, // 1-based Flash-X convention -> 0-based
gridIdxOrBlkId_{tI.gridIdxOrBlkId},
tileIdx_{tI.tileIdx},
lo_{LIST_NDIM(tI.loX,tI.loY,tI.loZ)},
hi_{LIST_NDIM(tI.hiX,tI.hiY,tI.hiZ)},
loGC_{LIST_NDIM(tI.loGCX,tI.loGCY,tI.loGCZ)},
hiGC_{LIST_NDIM(tI.hiGCX,tI.hiGCY,tI.hiGCZ)},
lo_{LIST_NDIM(tI.loX-1,tI.loY-1,tI.loZ-1)}, // 1-based Flash-X convention -> 0-based
hi_{LIST_NDIM(tI.hiX-1,tI.hiY-1,tI.hiZ-1)}, // 1-based Flash-X convention -> 0-based
loGC_{LIST_NDIM(tI.loGCX-1,tI.loGCY-1,tI.loGCZ-1)}, // 1-based Flash-X convention -> 0-based
hiGC_{LIST_NDIM(tI.hiGCX-1,tI.hiGCY-1,tI.hiGCZ-1)}, // 1-based Flash-X convention -> 0-based
nCcComp_{tI.nCcComp},
nFluxComp_{tI.nFluxComp},
deltas_{LIST_NDIM(tR.deltaX,tR.deltaY,tR.deltaZ)},
Expand Down
19 changes: 11 additions & 8 deletions tools/milhoja_pypkg/src/milhoja/FortranTemplateUtility.py
Original file line number Diff line number Diff line change
Expand Up @@ -168,6 +168,7 @@ def iterate_tilemetadata(

construct_host = ""
if pure_source != LBOUND_ARGUMENT:
# adjusting index base
fix_index = '+1' if pure_source in bounds_data else ''
info.dtype = F2C_TYPE_MAPPING.get(info.dtype, info.dtype)

Expand All @@ -176,15 +177,17 @@ def iterate_tilemetadata(
# not have knowledge of the other
if source == TILE_INTERIOR_ARGUMENT:
construct_host = "[MILHOJA_MDIM * 2] = {" \
"tileDesc_h->lo().I()+1,tileDesc_h->hi().I()+1, " \
"tileDesc_h->lo().J()+1,tileDesc_h->hi().J()+1, " \
"tileDesc_h->lo().K()+1,tileDesc_h->hi().K()+1 }"
f"tileDesc_h->lo().I(){fix_index}, tileDesc_h->hi().I(){fix_index}, " \
f"tileDesc_h->lo().J(){fix_index}, tileDesc_h->hi().J(){fix_index}, " \
f"tileDesc_h->lo().K(){fix_index}, tileDesc_h->hi().K(){fix_index} " \
"}"

elif source == TILE_ARRAY_BOUNDS_ARGUMENT:
construct_host = "[MILHOJA_MDIM * 2] = {" \
"tileDesc_h->loGC().I()+1,tileDesc_h->hiGC().I()+1, "\
"tileDesc_h->loGC().J()+1,tileDesc_h->hiGC().J()+1, "\
"tileDesc_h->loGC().K()+1,tileDesc_h->hiGC().K()+1 }"
f"tileDesc_h->loGC().I(){fix_index}, tileDesc_h->hiGC().I(){fix_index}, "\
f"tileDesc_h->loGC().J(){fix_index}, tileDesc_h->hiGC().J(){fix_index}, "\
f"tileDesc_h->loGC().K(){fix_index}, tileDesc_h->hiGC().K(){fix_index} " \
"}"

elif source == TILE_LEVEL_ARGUMENT:
one_time_mdata[item] = data
Expand All @@ -202,8 +205,8 @@ def iterate_tilemetadata(
else:
# info.dtype = VECTOR_ARRAY_EQUIVALENT[info.dtype]
# intvect i,j,k start at 0 so we need to add 1 to the
# index. However, anything that's just integers needs to be
# untouched.
# index as it is a Fortran.
# However, anything that's just integers needs to be untouched.
for idx, value in enumerate(lbound):
found = re.search('[a-zA-z]', value)
if found:
Expand Down
44 changes: 23 additions & 21 deletions tools/milhoja_pypkg/src/milhoja/TaskFunctionCpp2CGenerator_cpu_F.py
Original file line number Diff line number Diff line change
Expand Up @@ -16,7 +16,8 @@
TILE_UBOUND_ARGUMENT, SCRATCH_ARGUMENT, F2C_TYPE_MAPPING,
THREAD_INDEX_VAR_NAME, GRID_DATA_ARGUMENT, TILE_INTERIOR_ARGUMENT,
TILE_ARRAY_BOUNDS_ARGUMENT, GRID_DATA_PTRS, SOURCE_DATATYPES,
VECTOR_ARRAY_EQUIVALENT, TILE_ARGUMENTS_ALL, GRID_DATA_LBOUNDS
VECTOR_ARRAY_EQUIVALENT, TILE_ARGUMENTS_ALL, GRID_DATA_LBOUNDS,
TILE_INDEX_DATA,
)


Expand Down Expand Up @@ -149,21 +150,18 @@ def _fill_mdata_connectors(self, arg, spec, connectors: dict, saved):
if src == TILE_INTERIOR_ARGUMENT or src == TILE_ARRAY_BOUNDS_ARGUMENT:
lo = 'tile_lo' if src == TILE_INTERIOR_ARGUMENT else "tile_loGC"
hi = 'tile_hi' if src == TILE_INTERIOR_ARGUMENT else "tile_hiGC"
# Fortran assumes 1-based index
offs = "+1"
connectors[self.C2F_ARG_LIST].append(f"const void* {src}")
connectors[self.REAL_ARGS].append(f"static_cast<void*>({src})")
lo_data = lo.replace("tile_", "") + "()"
hi_data = hi.replace("tile_", "") + "()"
lo_fmt = f"{self.tile_desc_name}->{lo_data}." + '{0}'
hi_fmt = f"{self.tile_desc_name}->{hi_data}." + '{0}'
combined = f"int {src}[] = {{\n{self.INDENT}"
combined += f',\n{self.INDENT}'.join(
'{0}, {1}'.format(
indexexp_fmt.format(lo_fmt.format(char)),
indexexp_fmt.format(hi_fmt.format(char))
'{0}->{1}.{3}(){4}, {0}->{2}.{3}(){4}'.format(
self.tile_desc_name, lo_data, hi_data, char, offs
)
for indexexp_fmt, char in [('{0}', 'I()'),
('IFELSE_K2D({0},1)', 'J()'),
('IFELSE_K3D({0},1)', 'K()')]
for char in ['I', 'J', 'K']
) + "\n}"

self.connectors[self.CONSOLIDATE_TILE_DATA].append(combined)
Expand Down Expand Up @@ -191,19 +189,16 @@ def _fill_mdata_connectors(self, arg, spec, connectors: dict, saved):

dtype = SOURCE_DATATYPES[src]
if dtype in VECTOR_ARRAY_EQUIVALENT:
offs = ""
if src in TILE_INDEX_DATA:
# Fortran assumes 1-based index
offs = "+1"
raw = VECTOR_ARRAY_EQUIVALENT[dtype]
if raw == "int":
connectors[self.CONSOLIDATE_TILE_DATA].append(
f"{raw} {arg}_array[] = {{\n{self.INDENT}{arg}.I(),\n"
f"{self.INDENT}IFELSE_K2D({arg}.J(),1),\n"
f"{self.INDENT}IFELSE_K3D({arg}.K(),1)\n}}"
)
else: # for arguments like tile_deltas
connectors[self.CONSOLIDATE_TILE_DATA].append(
f"{raw} {arg}_array[] = {{\n{self.INDENT}{arg}.I(),\n"
f"{self.INDENT}{arg}.J(),\n"
f"{self.INDENT}{arg}.K()\n}}"
)
connectors[self.CONSOLIDATE_TILE_DATA].append(
f"{raw} {arg}_array[] = {{\n{self.INDENT}{arg}.I(){offs},\n"
f"{self.INDENT}{arg}.J(){offs},\n"
f"{self.INDENT}{arg}.K(){offs}\n}}"
)

else:
if 'unsigned' in dtype:
Expand Down Expand Up @@ -253,6 +248,13 @@ def _fill_lbound_connectors(self, arg, spec, connectors, saved):
)
saved.add(word)

# adjusting the base index
# because it is a Fortran
for i, bound in enumerate(lb):
for keyword in TILE_INDEX_DATA:
if keyword in bound:
lb[i] = bound + "+1"

lb = f"{{\n{self.INDENT}" + f',\n{self.INDENT}'.join(lb) + "\n}"
lb = lb.replace(" ", "")
connectors[self.CONSOLIDATE_TILE_DATA].append(
Expand Down
1 change: 1 addition & 0 deletions tools/milhoja_pypkg/src/milhoja/__init__.py
Original file line number Diff line number Diff line change
Expand Up @@ -34,6 +34,7 @@
TILE_FACE_AREAS_ARGUMENT,
TILE_CELL_VOLUMES_ARGUMENT,
TILE_ARGUMENTS_ALL,
TILE_INDEX_DATA,
GRID_DATA_LBOUNDS, GRID_DATA_EXTENTS,
C2F_TYPE_MAPPING, VECTOR_ARRAY_EQUIVALENT,
GRID_DATA_PTRS, SOURCE_DATATYPES, F2C_TYPE_MAPPING,
Expand Down
10 changes: 10 additions & 0 deletions tools/milhoja_pypkg/src/milhoja/constants.py
Original file line number Diff line number Diff line change
Expand Up @@ -66,6 +66,16 @@
TILE_CELL_VOLUMES_ARGUMENT
}

# The following index data must be handled carefully
# with a proper starting index, e.g., 1-based for Fortran language,
# during the code generations.
TILE_INDEX_DATA = [
TILE_LO_ARGUMENT,
TILE_HI_ARGUMENT,
TILE_LBOUND_ARGUMENT,
TILE_UBOUND_ARGUMENT,
]

# ----- DATA ACCESS KEYWORDS
READ = "r"
WRITE = "w"
Expand Down
21 changes: 5 additions & 16 deletions tools/milhoja_pypkg/src/milhoja/parse_helpers.py
Original file line number Diff line number Diff line change
Expand Up @@ -218,7 +218,7 @@ def parse_lbound_f(lbound: str):
)

# find everything between a single set of parens to find math symbols.
regexr = r'\(([^\(\)]+|(?:[^\(\)]|\(\))+)\)'
regexr = r'\(([^\)]+)\)'
matches = re.findall(regexr, lbound)
math_sym_string = lbound
for match in matches:
Expand All @@ -230,30 +230,19 @@ def parse_lbound_f(lbound: str):
for idx, match in enumerate(matches):
for keyword in keywords:
if keyword in match:
if (("IFELSE_K2D(" in match) or ("IFELSE_K3D(" in match)):
matches[idx] = match.replace(
keyword, f'{keyword}.I(),{keyword}.J(),{keyword}.K()'
)
else:
matches[idx] = match.replace(
keyword, f'{keyword}.I(),IFELSE_K2D({keyword}.J();1),IFELSE_K3D({keyword}.K();1)'
)
matches[idx] = match.replace(
keyword, f'{keyword}.I(),{keyword}.J(),{keyword}.K()'
)

iterables = [match.split(',') for match in matches]
if not iterables:
raise RuntimeError(f"Nothing in lbound {lbound}.")

# Replace ';' characters, which were temporarily used in place of ',' in the string
# in order to prevent splitting there, by the intended ';'.
for idx, iterable in enumerate(iterables):
for id, match in enumerate(iterables[idx]):
iterables[idx][id] = match.replace("();1)", "(),1)")

# check if all lists inside the bounds equations are the same length.
# Don't attempt to stitch different length arrays together.
size = len(iterables[0])
if not all([len(item) == size for item in iterables]):
raise RuntimeError(f"Different lbound part sizes for {lbound}: {iterables}")
raise RuntimeError(f"Different lbound part sizes. {lbound}")

# combine all lbound parts into 1.
# list of mathematic expressions will always be 1 less than the number
Expand Down
8 changes: 4 additions & 4 deletions tools/milhoja_pypkg/src/milhoja/tests/TestParseHelpers.py
Original file line number Diff line number Diff line change
Expand Up @@ -96,7 +96,7 @@ def check_bound(self, input, generated, correct):
"""
self.assertTrue(
generated == correct,
f"{input} returned {generated}, instead of {correct}."
f"{input} returned {generated}, istead of {correct}."
)

def testLboundParser(self):
Expand Down Expand Up @@ -206,7 +206,7 @@ def testLboundFParser(self):
* Test found keywords
"""
lb_input = "(tile_lo) - (1,2,3) + (2,3,4)"
correct = ["tile_lo.I()-1+2", "IFELSE_K2D(tile_lo.J(),1)-2+3", "IFELSE_K3D(tile_lo.K(),1)-3+4"]
correct = ["tile_lo.I()-1+2", "tile_lo.J()-2+3", "tile_lo.K()-3+4"]
result, _ = parse_lbound_f(lb_input)
self.check_bound(lb_input, result, correct)

Expand All @@ -218,7 +218,7 @@ def testLboundFParser(self):
result, _ = parse_lbound_f(lb_input)

lb_input = "(tile_hi, -1)"
correct = ["tile_hi.I()", "IFELSE_K2D(tile_hi.J(),1)", "IFELSE_K3D(tile_hi.K(),1)", "-1"]
correct = ["tile_hi.I()", "tile_hi.J()", "tile_hi.K()", "-1"]
result, _ = parse_lbound_f(lb_input)
self.check_bound(lb_input, result, correct)

Expand All @@ -228,7 +228,7 @@ def testLboundFParser(self):
self.check_bound(lb_input, result, correct)

lb_input = "(1, tile_lo, 6)"
correct = ["1", "tile_lo.I()", "IFELSE_K2D(tile_lo.J(),1)", "IFELSE_K3D(tile_lo.K(),1)", '6']
correct = ["1", "tile_lo.I()", "tile_lo.J()", "tile_lo.K()", '6']
result, _ = parse_lbound_f(lb_input)
self.check_bound(lb_input, result, correct)

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -293,11 +293,11 @@ void DataPacket_gpu_tf_hydro::pack(void) {
char_ptr = static_cast<char*>(static_cast<void*>(_tile_arrayBounds_p)) + n * SIZE_TILE_ARRAYBOUNDS;
std::memcpy(static_cast<void*>(char_ptr), static_cast<void*>(_tile_arrayBounds_h), SIZE_TILE_ARRAYBOUNDS);

int _lbdd_CC_1_h[4] = {(lbound.I()) + 1,(IFELSE_K2D(lbound.J(),1)) + 1,(IFELSE_K3D(lbound.K(),1)) + 1,1};
int _lbdd_CC_1_h[4] = {(lbound.I()) + 1,(lbound.J()) + 1,(lbound.K()) + 1,1};
char_ptr = static_cast<char*>(static_cast<void*>(_lbdd_CC_1_p)) + n * SIZE_LBDD_CC_1;
std::memcpy(static_cast<void*>(char_ptr), static_cast<void*>(_lbdd_CC_1_h), SIZE_LBDD_CC_1);

int _lbdd_scratch_hydro_op1_auxC_h[3] = {(lo.I()-1) + 1,(IFELSE_K2D(lo.J(),1)- 1) + 1,(IFELSE_K3D(lo.K(),1)- 1) + 1};
int _lbdd_scratch_hydro_op1_auxC_h[3] = {(lo.I()-1) + 1,(lo.J()- 1) + 1,(lo.K()- 1) + 1};
char_ptr = static_cast<char*>(static_cast<void*>(_lbdd_scratch_hydro_op1_auxC_p)) + n * SIZE_LBDD_SCRATCH_HYDRO_OP1_AUXC;
std::memcpy(static_cast<void*>(char_ptr), static_cast<void*>(_lbdd_scratch_hydro_op1_auxC_h), SIZE_LBDD_SCRATCH_HYDRO_OP1_AUXC);

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -106,40 +106,40 @@ extern "C" {
};
int tile_hi_array[] = {
tile_hi.I(),
IFELSE_K2D(tile_hi.J(),1),
IFELSE_K3D(tile_hi.K(),1)
tile_hi.J(),
tile_hi.K()
};
int tile_interior[] = {
tileDesc->lo().I(), tileDesc->hi().I(),
IFELSE_K2D(tileDesc->lo().J(),1), IFELSE_K2D(tileDesc->hi().J(),1),
IFELSE_K3D(tileDesc->lo().K(),1), IFELSE_K3D(tileDesc->hi().K(),1)
tileDesc->lo().I(),tileDesc->hi().I(),
tileDesc->lo().J(),tileDesc->hi().J(),
tileDesc->lo().K(),tileDesc->hi().K()
};
int tile_lo_array[] = {
tile_lo.I(),
IFELSE_K2D(tile_lo.J(),1),
IFELSE_K3D(tile_lo.K(),1)
tile_lo.J(),
tile_lo.K()
};
int lbdd_CC_1[] = {
tile_lbound.I(),
IFELSE_K2D(tile_lbound.J(),1),
IFELSE_K3D(tile_lbound.K(),1),
tile_lbound.J(),
tile_lbound.K(),
1
};
int lbdd_scratch_hydro_op1_auxC[] = {
tile_lo.I()-1,
IFELSE_K2D(tile_lo.J(),1)-1,
IFELSE_K3D(tile_lo.K(),1)-0
tile_lo.J()-1,
tile_lo.K()-0
};
int lbdd_scratch_hydro_op1_flX[] = {
tile_lo.I(),
IFELSE_K2D(tile_lo.J(),1),
IFELSE_K3D(tile_lo.K(),1),
tile_lo.J(),
tile_lo.K(),
1
};
int lbdd_scratch_hydro_op1_flY[] = {
tile_lo.I(),
IFELSE_K2D(tile_lo.J(),1),
IFELSE_K3D(tile_lo.K(),1),
tile_lo.J(),
tile_lo.K(),
1
};

Expand Down
Loading

0 comments on commit 6c1e22a

Please sign in to comment.