Skip to content

resolve issues around MAX_CONST_SIZE to address #253 #254

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

Merged
bozbez merged 1 commit into from
Nov 6, 2025
Merged

resolve issues around MAX_CONST_SIZE to address #253 #254

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
5 changes: 0 additions & 5 deletions translator-v2/resources/templates/cpp/master_kernel.cpp.jinja
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,11 +20,6 @@ extern {{const.typ}} {{const.ptr}}{% if const.dim > 1 %}[{{const.dim}}]{% endif
void op_decl_const_char(int dim, const char *type, int size, char *dat, const char *name) {
{{guard}}

if (size > MAX_CONST_SIZE) {
printf("error: requested size %d for const %s exceeds MAX_CONST_SIZE\n", size, name);
exit(1);
}

{% for const in app.consts() %}
if (!strcmp(name, "{{const.ptr}}")) {
{{caller(const)}}
Expand Down
14 changes: 7 additions & 7 deletions translator/c/op2.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -90,7 +90,7 @@ def op_decl_set_parse(text):
"""Parsing for op_decl_set calls"""

sets = []
for m in re.finditer('op_decl_set\((.*)\)', text):
for m in re.finditer(r'op_decl_set\((.*)\)', text):
args = m.group(1).split(',')

# check for syntax errors
Expand All @@ -101,7 +101,7 @@ def op_decl_set_parse(text):
sets.append({
'name': args[1].strip()
})
for m in re.finditer('op_decl_set_hdf5\((.*)\)', text):
for m in re.finditer(r'op_decl_set_hdf5\((.*)\)', text):
args = m.group(1).split(',')

# check for syntax errors
Expand All @@ -120,7 +120,7 @@ def op_decl_const_parse(text):
"""Parsing for op_decl_const calls"""

consts = []
for m in re.finditer('op_decl_const\((.*)\)', text):
for m in re.finditer(r'op_decl_const\((.*)\)', text):
args = m.group(1).split(',')

# check for syntax errors
Expand Down Expand Up @@ -242,8 +242,8 @@ def get_arg_gbl(arg_string, k):
return temp_gbl

def append_init_soa(text):
text = re.sub('\\bop_init\\b\\s*\((.*)\)','op_init_soa(\\1,1)', text)
text = re.sub('\\bop_mpi_init\\b\\s*\((.*)\)','op_mpi_init_soa(\\1,1)', text)
text = re.sub(r'\\bop_init\\b\\s*\((.*)\)','op_init_soa(\\1,1)', text)
text = re.sub(r'\\bop_mpi_init\\b\\s*\((.*)\)','op_mpi_init_soa(\\1,1)', text)
return text

def op_par_loop_parse(text):
Expand Down Expand Up @@ -305,9 +305,9 @@ def op_par_loop_parse(text):

def op_check_kernel_in_text(text, name):
match = False
inline_impl_pattern = r'inline[ \n]+void[ \n]+'+name+'\s*\('
inline_impl_pattern = r'inline[ \n]+void[ \n]+'+name+r'\s*\('
matches = re.findall(inline_impl_pattern, text)
decl_pattern = r'([$\n]+)(void[ \n]+'+name+'\([ \n]*'+'[ \nA-Za-z0-9\*\_\.,#]+\);)'
decl_pattern = r'([$\n]+)(void[ \n]+'+name+r'\([ \n]*'+r'[ \nA-Za-z0-9\*\_\.,#]+\);)'
if len(re.findall(inline_impl_pattern, text)) == 1:
match = True
elif len(re.findall(decl_pattern, text)) == 1:
Expand Down
8 changes: 4 additions & 4 deletions translator/c/op2_gen_cuda.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -961,10 +961,10 @@ def op2_gen_cuda(master, date, consts, kernels, sets):
code('')

for nc in range (0,len(consts)):
if consts[nc]['dim']==1:
if str(consts[nc]['dim']).isdigit() and int(consts[nc]['dim'])==1:
code('__constant__ '+consts[nc]['type'][1:-1]+' '+consts[nc]['name']+';')
else:
if consts[nc]['dim'] > 0:
if str(consts[nc]['dim']).isdigit() and int(consts[nc]['dim']) > 0:
num = str(consts[nc]['dim'])
else:
num = 'MAX_CONST_SIZE'
Expand Down Expand Up @@ -992,8 +992,8 @@ def op2_gen_cuda(master, date, consts, kernels, sets):
code(' '+consts[nc]['type'][1:-1]+' *dat){')
depth = depth + 2
code('if (!OP_hybrid_gpu) return;')
if not consts[nc]['dim'] or int(consts[nc]['dim']) > 1:
IF('dim*sizeof('+consts[nc]['type'][1:-1]+')>MAX_CONST_SIZE')
if not str(consts[nc]['dim']).isdigit() or int(consts[nc]['dim']) <= 0:
IF('dim>MAX_CONST_SIZE')
code('printf("error: MAX_CONST_SIZE not big enough\\n"); exit(1);')
ENDIF()
code('cutilSafeCall(cudaMemcpyToSymbol('+consts[nc]['name']+'_cuda, dat, dim*sizeof('+consts[nc]['type'][1:-1]+')));')
Expand Down
8 changes: 4 additions & 4 deletions translator/c/op2_gen_cuda_simple.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1291,10 +1291,10 @@ def op2_gen_cuda_simple(master, date, consts, kernels,sets, macro_defs):
code('')

for nc in range (0,len(consts)):
if consts[nc]['dim']==1:
if str(consts[nc]['dim']).isdigit() and int(consts[nc]['dim'])==1:
code('__constant__ '+consts[nc]['type'][1:-1]+' '+consts[nc]['name']+'_cuda;')
else:
if consts[nc]['dim'].isdigit() and int(consts[nc]['dim']) > 0:
if str(consts[nc]['dim']).isdigit() and int(consts[nc]['dim']) > 0:
num = str(consts[nc]['dim'])
else:
num = 'MAX_CONST_SIZE'
Expand All @@ -1317,8 +1317,8 @@ def op2_gen_cuda_simple(master, date, consts, kernels,sets, macro_defs):
code(' '+consts[nc]['type'][1:-1]+' *dat){')
depth = depth + 2
code('if (!OP_hybrid_gpu) return;')
if not consts[nc]['dim'] or int(consts[nc]['dim']) > 1:
IF('dim*sizeof('+consts[nc]['type'][1:-1]+')>MAX_CONST_SIZE')
if not str(consts[nc]['dim']).isdigit() or int(consts[nc]['dim']) <= 0:
IF('dim>MAX_CONST_SIZE')
code('printf("error: MAX_CONST_SIZE not big enough\\n"); exit(1);')
ENDIF()
code('cutilSafeCall(cudaMemcpyToSymbol('+consts[nc]['name']+'_cuda, dat, dim*sizeof('+consts[nc]['type'][1:-1]+')));')
Expand Down