def test_profile_invokes_gocean1p0():
'''Check that an invoke is instrumented correctly
'''
Profiler.set_options([Profiler.INVOKES])
_, invoke = get_invoke("test11_different_iterates_over_one_invoke.f90",
"gocean1.0",
idx=0)
# Convert the invoke to code, and remove all new lines, to make
# regex matching easier
code = str(invoke.gen()).replace("\n", "")
# First a simple test that the nesting is correct - the
# profile regions include both loops. Note that indeed
# the function 'compute_cv_code' is in the module file
# kernel_ne_offset_mod.
correct_re = ("subroutine invoke.*"
"use profile_mod, only: ProfileData.*"
r"TYPE\(ProfileData\), save :: profile.*"
r"call ProfileStart\(\"kernel_ne_offset_mod\", "
r"\"compute_cv_code\", profile\).*"
"do j.*"
"do i.*"
"call.*"
"end.*"
"end.*"
r"call ProfileEnd\(profile\)")
assert re.search(correct_re, code, re.I) is not None
# Check that if gen() is called more than once the same profile
# variables and region names are created:
code_again = str(invoke.gen()).replace("\n", "")
assert code == code_again
# Test that two kernels in one invoke get instrumented correctly.
_, invoke = get_invoke("single_invoke_two_kernels.f90", "gocean1.0", 0)
# Convert the invoke to code, and remove all new lines, to make
# regex matching easier
code = str(invoke.gen()).replace("\n", "")
correct_re = ("subroutine invoke.*"
"use profile_mod, only: ProfileData.*"
r"TYPE\(ProfileData\), save :: profile.*"
r"call ProfileStart\(\"compute_cu_mod\", "
r"\"compute_cu_code\", profile\).*"
"do j.*"
"do i.*"
"call.*"
"end.*"
"end.*"
"do j.*"
"do i.*"
"call.*"
"end.*"
"end.*"
r"call ProfileEnd\(profile\)")
assert re.search(correct_re, code, re.I) is not None
Profiler.set_options(None)
def test_profile_kernels_dynamo0p3():
'''Check that all kernels are instrumented correctly in a
Dynamo 0.3 invoke.
'''
Profiler.set_options([Profiler.KERNELS])
_, invoke = get_invoke("1_single_invoke.f90", "dynamo0.3", idx=0)
# Convert the invoke to code, and remove all new lines, to make
# regex matching easier
code = str(invoke.gen()).replace("\n", "")
correct_re = ("subroutine invoke.*"
"use profile_mod, only: ProfileData, ProfileStart, "
"ProfileEnd.*"
r"TYPE\(ProfileData\), save :: profile.*"
r"call ProfileStart\(\"testkern\", \"testkern_code.*\", "
r"profile\).*"
"do cell.*"
"call.*"
"end.*"
r"call ProfileEnd\(profile\)")
assert re.search(correct_re, code, re.I) is not None
_, invoke = get_invoke("1.2_multi_invoke.f90", "dynamo0.3", idx=0)
# Convert the invoke to code, and remove all new lines, to make
# regex matching easier
code = str(invoke.gen()).replace("\n", "")
correct_re = ("subroutine invoke.*"
"use profile_mod, only: ProfileData, ProfileStart, "
"ProfileEnd.*"
r"TYPE\(ProfileData\), save :: profile.*"
r"TYPE\(ProfileData\), save :: profile.*"
r"call ProfileStart\(\"testkern\", \"testkern_code.*\", "
r"(?P<profile1>\w*)\).*"
"do cell.*"
"call.*"
"end.*"
r"call ProfileEnd\((?P=profile1)\).*"
r"call ProfileStart\(.*, (?P<profile2>\w*)\).*"
"do cell.*"
"call.*"
"end.*"
r"call ProfileEnd\((?P=profile2)\).*")
groups = re.search(correct_re, code, re.I)
assert groups is not None
# Check that the variables are different
assert groups.group(1) != groups.group(2)
Profiler.set_options(None)
def test_no_inline_before_trans(monkeypatch, tmpdir):
''' Check that we reject attempts to transform kernels that have been
marked for module in-lining. Issue #229. '''
# Even though this code will raise an exception at the end, it will
# still create an (empty) file 'testkern_any_space_2_0_mod.f90'.
# So change to tmpdir to avoid this.
old_pwd = tmpdir.chdir()
from psyclone.transformations import KernelModuleInlineTrans
psy, invoke = get_invoke("4.5.2_multikernel_invokes.f90", api="dynamo0.3",
idx=0)
sched = invoke.schedule
kernels = sched.walk(Kern)
assert len(kernels) == 5
inline_trans = KernelModuleInlineTrans()
rtrans = ACCRoutineTrans()
_, _ = inline_trans.apply(kernels[1])
with pytest.raises(TransformationError) as err:
_, _ = rtrans.apply(kernels[1])
assert "because it will be module-inlined" in str(err)
# Monkeypatch the validate() routine so we can check that we catch
# the error at code-generation time
monkeypatch.setattr(rtrans, "validate", lambda kern: None)
_, _ = rtrans.apply(kernels[1])
with pytest.raises(NotImplementedError) as err:
_ = str(psy.gen).lower()
assert "Cannot module-inline a transformed kernel " in str(err)
old_pwd.chdir()
def test_new_same_kern_single(tmpdir, monkeypatch):
''' Check that we do not overwrite an existing, identical kernel if
there is a name clash and kernel-naming is 'single'. '''
# Ensure kernel-output directory is uninitialised
config = Config.get()
monkeypatch.setattr(config, "_kernel_output_dir", "")
monkeypatch.setattr(config, "_kernel_naming", "single")
rtrans = ACCRoutineTrans()
# Change to temp dir (so kernel written there)
old_cwd = tmpdir.chdir()
_, invoke = get_invoke("4_multikernel_invokes.f90", api="dynamo0.3",
idx=0)
sched = invoke.schedule
# Apply the same transformation to both kernels. This should produce
# two, identical transformed kernels.
new_kernels = []
for kern in sched.coded_kernels():
new_kern, _ = rtrans.apply(kern)
new_kernels.append(new_kern)
# Generate the code - we should end up with just one transformed kernel
new_kernels[0].rename_and_write()
new_kernels[1].rename_and_write()
assert new_kernels[1]._name == "testkern_0_code"
assert new_kernels[1].module_name == "testkern_0_mod"
out_files = os.listdir(str(tmpdir))
assert out_files == [new_kernels[1].module_name+".f90"]
old_cwd.chdir()
def test_new_kern_single_error(tmpdir, monkeypatch):
''' Check that we do not overwrite an existing, different kernel if
there is a name clash and kernel-naming is 'single'. '''
# Ensure kernel-output directory is uninitialised
config = Config.get()
monkeypatch.setattr(config, "_kernel_output_dir", "")
monkeypatch.setattr(config, "_kernel_naming", "single")
# Change to temp dir (so kernel written there)
old_cwd = tmpdir.chdir()
_, invoke = get_invoke("1_single_invoke.f90", api="dynamo0.3", idx=0)
sched = invoke.schedule
kernels = sched.coded_kernels()
kern = kernels[0]
old_mod_name = kern.module_name[:]
# Create a file with the same name as we would otherwise generate
with open(os.path.join(str(tmpdir),
old_mod_name+"_0_mod.f90"), "w") as ffile:
ffile.write("some code")
rtrans = ACCRoutineTrans()
new_kern, _ = rtrans.apply(kern)
# Generate the code - this should raise an error as we get a name
# clash and the content of the existing file is not the same as that
# which we would generate
with pytest.raises(GenerationError) as err:
new_kern.rename_and_write()
assert ("transformed version of this Kernel 'testkern_0_mod.f90' already "
"exists in the kernel-output directory ({0}) but is not the same "
"as the current, transformed kernel and the kernel-renaming "
"scheme is set to 'single'".format(str(tmpdir)) in str(err))
old_cwd.chdir()
def test_new_kern_no_clobber(tmpdir, monkeypatch):
''' Check that we create a new kernel with a new name when kernel-naming
is set to 'multiple' and we would otherwise get a name clash. '''
# Ensure kernel-output directory is uninitialised
config = Config.get()
monkeypatch.setattr(config, "_kernel_output_dir", "")
monkeypatch.setattr(config, "_kernel_naming", "multiple")
# Change to temp dir (so kernel written there)
old_cwd = tmpdir.chdir()
psy, invoke = get_invoke("1_single_invoke.f90", api="dynamo0.3", idx=0)
sched = invoke.schedule
kernels = sched.walk(Kern)
kern = kernels[0]
old_mod_name = kern.module_name[:]
# Create a file with the same name as we would otherwise generate
with open(os.path.join(str(tmpdir),
old_mod_name+"_0_mod.f90"), "w") as ffile:
ffile.write("some code")
rtrans = ACCRoutineTrans()
_, _ = rtrans.apply(kern)
# Generate the code (this triggers the generation of a new kernel)
_ = str(psy.gen).lower()
filename = os.path.join(str(tmpdir), old_mod_name+"_1_mod.f90")
assert os.path.isfile(filename)
old_cwd.chdir()
def test_accroutine_err(monkeypatch):
''' Check that we raise the expected error if we can't find the
source of the kernel subroutine. '''
import fparser
_, invoke = get_invoke("1_single_invoke.f90", api="dynamo0.3", idx=0)
sched = invoke.schedule
kernels = sched.coded_kernels()
kern = kernels[0]
assert isinstance(kern, Kern)
# Edit the fparser1 AST of the kernel so that it does not have a
# subroutine of the correct name
ast = kern._module_code
mod = ast.content[0]
# Find the subroutine statement
for child in mod.content:
if isinstance(child, fparser.one.block_statements.Subroutine):
sub = child
# Find the end subroutine statement
for child in sub.content:
if isinstance(child, fparser.one.block_statements.EndSubroutine):
end = child
monkeypatch.setattr(sub, "name", "some_other_name")
monkeypatch.setattr(end, "name", "some_other_name")
rtrans = ACCRoutineTrans()
with pytest.raises(TransformationError) as err:
_ = rtrans.apply(kern)
assert ("Failed to find subroutine source for kernel testkern_code"
in str(err))
def test_accroutine():
''' Test that we can transform a kernel by adding a "!$acc routine"
directive to it. '''
from psyclone.gocean1p0 import GOKern
from fparser.two import Fortran2003
_, invoke = get_invoke("nemolite2d_alg_mod.f90", api="gocean1.0", idx=0)
sched = invoke.schedule
kern = sched.children[0].loop_body[0].loop_body[0]
assert isinstance(kern, GOKern)
rtrans = ACCRoutineTrans()
assert rtrans.name == "ACCRoutineTrans"
new_kern, _ = rtrans.apply(kern)
# The transformation should have populated the fparser2 AST of
# the kernel...
assert new_kern._fp2_ast
assert isinstance(new_kern._fp2_ast, Fortran2003.Program)
# Check AST contains directive
comments = walk_ast(new_kern._fp2_ast.content, [Fortran2003.Comment])
assert len(comments) == 1
assert str(comments[0]) == "!$acc routine"
# Check that directive is in correct place (end of declarations)
gen = str(new_kern._fp2_ast)
assert ("REAL(KIND = go_wp), DIMENSION(:, :), INTENT(IN) :: sshn, sshn_u, "
"sshn_v, hu, hv, un, vn\n"
" !$acc routine\n"
" ssha (ji, jj) = 0.0_go_wp\n" in gen)
def test_psy_init(kernel_outputdir):
''' Check that we create a psy_init() routine that sets-up the
OpenCL environment. '''
psy, _ = get_invoke("single_invoke.f90", API, idx=0)
sched = psy.invokes.invoke_list[0].schedule
otrans = OCLTrans()
otrans.apply(sched)
generated_code = str(psy.gen)
expected = (
" SUBROUTINE psy_init()\n"
" USE fortcl, ONLY: ocl_env_init, add_kernels\n"
" CHARACTER(LEN=30) kernel_names(1)\n"
" LOGICAL, save :: initialised=.False.\n"
" ! Check to make sure we only execute this routine once\n"
" IF (.not. initialised) THEN\n"
" initialised = .True.\n"
" ! Initialise the OpenCL environment/device\n"
" CALL ocl_env_init\n"
" ! The kernels this PSy layer module requires\n"
" kernel_names(1) = \"compute_cu_code\"\n"
" ! Create the OpenCL kernel objects. Expects to find all of "
"the compiled\n"
" ! kernels in PSYCLONE_KERNELS_FILE.\n"
" CALL add_kernels(1, kernel_names)\n"
" END IF \n"
" END SUBROUTINE psy_init\n")
assert expected in generated_code
assert GOcean1p0OpenCLBuild(kernel_outputdir).code_compiles(psy)
def test_2kern_trans(kernel_outputdir):
''' Check that we generate correct code when we transform two kernels
within a single invoke. '''
psy, invoke = get_invoke("4.5.2_multikernel_invokes.f90",
api="dynamo0.3",
idx=0)
sched = invoke.schedule
kernels = sched.walk(Kern)
assert len(kernels) == 5
ktrans = Dynamo0p3KernelConstTrans()
_, _ = ktrans.apply(kernels[1], number_of_layers=100)
_, _ = ktrans.apply(kernels[2], number_of_layers=100)
# Generate the code (this triggers the generation of new kernels)
code = str(psy.gen).lower()
# Find the tags added to the kernel/module names
for match in re.finditer('use testkern_any_space_2(.+?)_mod', code):
tag = match.group(1)
assert ("use testkern_any_space_2{0}_mod, only: "
"testkern_any_space_2{0}_code".format(tag) in code)
assert "call testkern_any_space_2{0}_code(".format(tag) in code
filepath = os.path.join(str(kernel_outputdir),
"testkern_any_space_2{0}_mod.f90".format(tag))
assert os.path.isfile(filepath)
assert "nlayers = 100" in open(filepath).read()
assert "use testkern_any_space_2_mod, only" not in code
assert "call testkern_any_space_2_code(" not in code
assert Dynamo0p3Build(kernel_outputdir).code_compiles(psy)
def test_1kern_trans(kernel_outputdir):
''' Check that we generate the correct code when an invoke contains
the same kernel more than once but only one of them is transformed. '''
psy, invoke = get_invoke("4_multikernel_invokes.f90",
api="dynamo0.3",
idx=0)
sched = invoke.schedule
kernels = sched.coded_kernels()
# We will transform the second kernel but not the first
kern = kernels[1]
rtrans = ACCRoutineTrans()
_, _ = rtrans.apply(kern)
# Generate the code (this triggers the generation of a new kernel)
code = str(psy.gen).lower()
tag = re.search('use testkern(.+?)_mod', code).group(1)
# We should have a USE for the original kernel and a USE for the new one
assert "use testkern{0}_mod, only: testkern{0}_code".format(tag) in code
assert "use testkern, only: testkern_code" in code
# Similarly, we should have calls to both the original and new kernels
assert "call testkern_code(" in code
assert "call testkern{0}_code(".format(tag) in code
first = code.find("call testkern_code(")
second = code.find("call testkern{0}_code(".format(tag))
assert first < second
assert Dynamo0p3Build(kernel_outputdir).code_compiles(psy)
def test_set_kern_float_arg():
''' Check that we generate correct code to set a real, scalar kernel
argument. '''
psy, _ = get_invoke("single_invoke_scalar_float_arg.f90", API, idx=0)
sched = psy.invokes.invoke_list[0].schedule
otrans = OCLTrans()
otrans.apply(sched)
generated_code = str(psy.gen)
expected = '''\
SUBROUTINE bc_ssh_code_set_args(kernel_obj, nx, a_scalar, ssh_fld, ''' + \
'''xstop, tmask)
USE clfortran, ONLY: clSetKernelArg
USE iso_c_binding, ONLY: c_sizeof, c_loc, c_intptr_t
USE ocl_utils_mod, ONLY: check_status
REAL(KIND=go_wp), intent(in), target :: a_scalar
INTEGER ierr
INTEGER(KIND=c_intptr_t), target :: ssh_fld, xstop, tmask
INTEGER(KIND=c_intptr_t), target :: kernel_obj
INTEGER, target :: nx
'''
assert expected in generated_code
expected = '''\
! Set the arguments for the bc_ssh_code OpenCL Kernel
ierr = clSetKernelArg(kernel_obj, 0, C_SIZEOF(nx), C_LOC(nx))
ierr = clSetKernelArg(kernel_obj, 1, C_SIZEOF(a_scalar), C_LOC(a_scalar))
CALL check_status('clSetKernelArg: arg 1 of bc_ssh_code', ierr)
ierr = clSetKernelArg(kernel_obj, 2, C_SIZEOF(ssh_fld), C_LOC(ssh_fld))
CALL check_status('clSetKernelArg: arg 2 of bc_ssh_code', ierr)
ierr = clSetKernelArg(kernel_obj, 3, C_SIZEOF(xstop), C_LOC(xstop))
CALL check_status('clSetKernelArg: arg 3 of bc_ssh_code', ierr)
ierr = clSetKernelArg(kernel_obj, 4, C_SIZEOF(tmask), C_LOC(tmask))
CALL check_status('clSetKernelArg: arg 4 of bc_ssh_code', ierr)
END SUBROUTINE bc_ssh_code_set_args'''
assert expected in generated_code
def test_opencl_kernel_code_generation():
# pylint: disable=invalid-name
''' Tests that gen_ocl method of the GOcean Kernel Schedule generates
the expected OpenCL code.
'''
from psyclone.psyir.backend.opencl import OpenCLWriter
psy, _ = get_invoke("single_invoke.f90", API, idx=0)
sched = psy.invokes.invoke_list[0].schedule
kernel = sched.children[0].loop_body[0].loop_body[0] # compute_cu kernel
kschedule = kernel.get_kernel_schedule()
expected_code = ("__kernel void compute_cu_code(\n"
" __global double * restrict cu,\n"
" __global double * restrict p,\n"
" __global double * restrict u\n"
" ){\n"
" int cuLEN1 = get_global_size(0);\n"
" int cuLEN2 = get_global_size(1);\n"
" int pLEN1 = get_global_size(0);\n"
" int pLEN2 = get_global_size(1);\n"
" int uLEN1 = get_global_size(0);\n"
" int uLEN2 = get_global_size(1);\n"
" int i = get_global_id(0);\n"
" int j = get_global_id(1);\n"
" cu[j * cuLEN1 + i] = ((0.5e0 * (p[j * pLEN1 + (i + 1)]"
" + p[j * pLEN1 + i])) * u[j * uLEN1 + i]);\n"
"}\n")
openclwriter = OpenCLWriter()
assert expected_code == openclwriter(kschedule)
def test_2kern_trans(tmpdir, monkeypatch):
''' Check that we generate correct code when we transform two kernels
within a single invoke. '''
# Ensure kernel-output directory is uninitialised
config = Config.get()
monkeypatch.setattr(config, "_kernel_output_dir", "")
# Change to temp dir (so kernel written there)
old_cwd = tmpdir.chdir()
psy, invoke = get_invoke("4.5.2_multikernel_invokes.f90", api="dynamo0.3",
idx=0)
sched = invoke.schedule
kernels = sched.walk(Kern)
assert len(kernels) == 5
rtrans = ACCRoutineTrans()
_, _ = rtrans.apply(kernels[1])
_, _ = rtrans.apply(kernels[2])
# Generate the code (this triggers the generation of new kernels)
code = str(psy.gen).lower()
# Find the tags added to the kernel/module names
for match in re.finditer('use testkern_any_space_2(.+?)_mod', code):
tag = match.group(1)
assert ("use testkern_any_space_2{0}_mod, only: "
"testkern_any_space_2{0}_code".format(tag) in code)
assert "call testkern_any_space_2{0}_code(".format(tag) in code
assert os.path.isfile(
os.path.join(str(tmpdir),
"testkern_any_space_2{0}_mod.f90".format(tag)))
assert "use testkern_any_space_2_mod, only" not in code
assert "call testkern_any_space_2_code(" not in code
assert Dynamo0p3Build(tmpdir).code_compiles(psy)
old_cwd.chdir()
def test_1kern_trans(tmpdir, monkeypatch):
''' Check that we generate the correct code when an invoke contains
the same kernel more than once but only one of them is transformed. '''
# Ensure kernel-output directory is uninitialised
config = Config.get()
monkeypatch.setattr(config, "_kernel_output_dir", "")
# Change to temp dir (so kernel written there)
old_cwd = tmpdir.chdir()
psy, invoke = get_invoke("4_multikernel_invokes.f90", api="dynamo0.3",
idx=0)
sched = invoke.schedule
kernels = sched.coded_kernels()
# We will transform the second kernel but not the first
kern = kernels[1]
rtrans = ACCRoutineTrans()
_, _ = rtrans.apply(kern)
# Generate the code (this triggers the generation of a new kernel)
code = str(psy.gen).lower()
tag = re.search('use testkern(.+?)_mod', code).group(1)
# We should have a USE for the original kernel and a USE for the new one
assert "use testkern{0}_mod, only: testkern{0}_code".format(tag) in code
assert "use testkern, only: testkern_code" in code
# Similarly, we should have calls to both the original and new kernels
assert "call testkern_code(" in code
assert "call testkern{0}_code(".format(tag) in code
first = code.find("call testkern_code(")
second = code.find("call testkern{0}_code(".format(tag))
assert first < second
assert Dynamo0p3Build(tmpdir).code_compiles(psy)
old_cwd.chdir()
def test_profile_invokes_dynamo0p3():
'''Check that a Dynamo 0.3 invoke is instrumented correctly
'''
Profiler.set_options([Profiler.INVOKES])
# First test for a single invoke with a single kernel work as expected:
_, invoke = get_invoke("1_single_invoke.f90", "dynamo0.3", idx=0)
# Convert the invoke to code, and remove all new lines, to make
# regex matching easier
code = str(invoke.gen()).replace("\n", "")
correct_re = ("subroutine invoke.*"
"use profile_mod, only: ProfileData.*"
r"TYPE\(ProfileData\), save :: profile.*"
r"call ProfileStart\(\"testkern\", \"testkern_code\", "
r"profile\).*"
"do cell.*"
"call.*"
"end.*"
r"call ProfileEnd\(profile\)")
assert re.search(correct_re, code, re.I) is not None
# Next test two kernels in one invoke:
_, invoke = get_invoke("1.2_multi_invoke.f90", "dynamo0.3", idx=0)
# Convert the invoke to code, and remove all new lines, to make
# regex matching easier
code = str(invoke.gen()).replace("\n", "")
# The .* after testkern_code is necessary since the name can be changed
# by PSyclone to avoid name duplications.
correct_re = ("subroutine invoke.*"
"use profile_mod, only: ProfileData.*"
r"TYPE\(ProfileData\), save :: profile.*"
r"call ProfileStart\(\"testkern\", \"testkern_code.*\","
r" profile\).*"
"do cell.*"
"call.*"
"end.*"
"do cell.*"
"call.*"
"end.*"
r"call ProfileEnd\(profile\)")
assert re.search(correct_re, code, re.I) is not None
Profiler.set_options(None)
def test_omp_transform():
'''Tests that the profiling transform works correctly with OMP
parallelisation.'''
_, invoke = get_invoke("test27_loop_swap.f90",
"gocean1.0",
name="invoke_loop1")
schedule = invoke.schedule
prt = ProfileRegionTrans()
omp_loop = GOceanOMPLoopTrans()
omp_par = OMPParallelTrans()
# Parallelise the first loop:
sched1, _ = omp_loop.apply(schedule.children[0])
sched2, _ = omp_par.apply(sched1.children[0])
sched3, _ = prt.apply(sched2.children[0])
correct = (
" CALL ProfileStart(\"boundary_conditions_ne_offset_mod\", "
"\"bc_ssh_code\", profile)\n"
" !$omp parallel default(shared), private(i,j)\n"
" !$omp do schedule(static)\n"
" DO j=2,jstop\n"
" DO i=2,istop\n"
" CALL bc_ssh_code(i, j, 1, t%data, t%grid%tmask)\n"
" END DO \n"
" END DO \n"
" !$omp end do\n"
" !$omp end parallel\n"
" CALL ProfileEnd(profile)")
code = str(invoke.gen())
assert correct in code
# Now add another profile node between the omp parallel and omp do
# directives:
sched3, _ = prt.apply(sched3.children[0].children[0].children[0])
code = str(invoke.gen())
correct = ''' CALL ProfileStart("boundary_conditions_ne_offset_mod", \
"bc_ssh_code", profile)
!$omp parallel default(shared), private(i,j)
CALL ProfileStart("boundary_conditions_ne_offset_mod", "bc_ssh_code_1", \
profile_1)
!$omp do schedule(static)
DO j=2,jstop
DO i=2,istop
CALL bc_ssh_code(i, j, 1, t%data, t%grid%tmask)
END DO\x20
END DO\x20
!$omp end do
CALL ProfileEnd(profile_1)
!$omp end parallel
CALL ProfileEnd(profile)'''
assert correct in code
def test_accroutine_empty_kernel():
''' Check that the directive goes at the end of the declarations,
even when the rest of the kernel is empty. '''
_, invoke = get_invoke("1_single_invoke.f90", api="dynamo0.3", idx=0)
sched = invoke.schedule
kernels = sched.coded_kernels()
rtrans = ACCRoutineTrans()
new_kern, _ = rtrans.apply(kernels[0])
# Check that directive is in correct place (end of declarations)
gen = str(new_kern._fp2_ast).lower()
assert "!$acc routine\n end subroutine testkern_code" in gen
def test_set_arg_const_scalar():
''' Check that an invoke that passes a scalar kernel argument by
value is rejected. (We haven't yet implemented the necessary code for
setting the value of such an argument in OpenCL.) '''
psy, _ = get_invoke("test00.1_invoke_kernel_using_const_scalar.f90",
API, idx=0)
sched = psy.invokes.invoke_list[0].schedule
otrans = OCLTrans()
with pytest.raises(NotImplementedError) as err:
otrans.apply(sched)
assert ("Cannot generate OpenCL for Invokes that contain kernels with "
"arguments passed by value" in str(err))
def test_builtin_no_trans():
''' Check that we reject attempts to transform built-in kernels. '''
from psyclone.dynamo0p3_builtins import DynBuiltIn
_, invoke = get_invoke("15.1.1_X_plus_Y_builtin.f90",
api="dynamo0.3", idx=0)
sched = invoke.schedule
kernels = sched.walk(DynBuiltIn)
rtrans = ACCRoutineTrans()
with pytest.raises(TransformationError) as err:
_ = rtrans.apply(kernels[0])
assert ("ACCRoutineTrans to a built-in kernel is not yet supported and "
"kernel 'x_plus_y' is of type " in str(err))
def test_get_invoke():
'''Tests get_invokes. '''
# First test all 4 valid APIs - we only make sure that no exception
# is raised, so no assert required
_, _ = get_invoke("openmp_fuse_test.f90", "gocean0.1", idx=0)
_, _ = get_invoke("test14_module_inline_same_kernel.f90",
"gocean1.0",
idx=0)
_, _ = get_invoke("algorithm/1_single_function.f90", "dynamo0.1", idx=0)
_, _ = get_invoke("1_single_invoke.f90", "dynamo0.3", idx=0)
# Test that an invalid name raises an exception
with pytest.raises(RuntimeError) as excinfo:
_, _ = get_invoke("test11_different_iterates_over_one_invoke.f90",
"gocean1.0",
name="invalid_name")
assert "Cannot find an invoke named 'invalid_name'" in str(excinfo)
# Test that an invalid API raises the right exception:
with pytest.raises(RuntimeError) as excinfo:
_, _ = get_invoke("test11_different_iterates_over_one_invoke.f90",
"invalid-api",
name="invalid_name")
assert "The API 'invalid-api' is not supported" in str(excinfo)
def test_new_kernel_dir(kernel_outputdir):
''' Check that we write out the transformed kernel to a specified
directory. '''
psy, invoke = get_invoke("nemolite2d_alg_mod.f90", api="gocean1.0", idx=0)
sched = invoke.schedule
kern = sched.children[0].loop_body[0].loop_body[0]
rtrans = ACCRoutineTrans()
_, _ = rtrans.apply(kern)
# Generate the code (this triggers the generation of a new kernel)
_ = str(psy.gen)
file_list = os.listdir(str(kernel_outputdir))
assert len(file_list) == 1
assert file_list[0] == 'continuity_0_mod.f90'
def test_unique_region_names():
'''Test that unique region names are created even when the kernel
names are identical.'''
Profiler.set_options([Profiler.KERNELS])
_, invoke = get_invoke("single_invoke_two_identical_kernels.f90",
"gocean1.0", 0)
# Convert the invoke to code, and remove all new lines, to make
# regex matching easier
code = str(invoke.gen()).replace("\n", "")
# This regular expression puts the region names into groups.
# Make sure even though the kernels have the same name, that
# the created regions have different names. In order to be
# flexible for future changes, we get the region names from
# the ProfileStart calls using a regular expressions (\w*
# being the group name enclosed in "") group. Python will store
# those two groups and they can be accessed using the resulting
# re object.group(n).
correct_re = ("subroutine invoke.*"
"use profile_mod, only: ProfileData.*"
r"TYPE\(ProfileData\), save :: profile.*"
r"TYPE\(ProfileData\), save :: profile.*"
r"call ProfileStart\(\"compute_cu_mod\", \"(\w*)\", "
r"profile\).*"
"do j.*"
"do i.*"
"call compute_cu_code.*"
"end.*"
"end.*"
r"call ProfileEnd\(profile\).*"
r"call ProfileStart\(\"compute_cu_mod\", \"(\w*)\", "
r"profile_1\).*"
"do j.*"
"do i.*"
"call compute_cu_code.*"
"end.*"
"end.*"
r"call ProfileEnd\(profile_1\)")
groups = re.search(correct_re, code, re.I)
assert groups is not None
# Check that the region names are indeed different: group(1)
# is the first kernel region name crated by PSyclone, and
# group(2) the name used in the second ProfileStart.
# Those names must be different (otherwise the profiling tool
# would likely combine the two different regions into one).
assert groups.group(1) != groups.group(2)
def test_new_kernel_file(tmpdir, monkeypatch):
''' Check that we write out the transformed kernel to the CWD. '''
from fparser.two import Fortran2003, parser
from fparser.common.readfortran import FortranFileReader
# Ensure kernel-output directory is uninitialised
config = Config.get()
monkeypatch.setattr(config, "_kernel_output_dir", "")
monkeypatch.setattr(config, "_kernel_naming", "multiple")
# Change to temp dir (so kernel written there)
old_cwd = tmpdir.chdir()
psy, invoke = get_invoke("nemolite2d_alg_mod.f90", api="gocean1.0", idx=0)
sched = invoke.schedule
kern = sched.children[0].loop_body[0].loop_body[0]
rtrans = ACCRoutineTrans()
_, _ = rtrans.apply(kern)
# Generate the code (this triggers the generation of a new kernel)
code = str(psy.gen).lower()
# Work out the value of the tag used to re-name the kernel
tag = re.search('use continuity(.+?)_mod', code).group(1)
assert ("use continuity{0}_mod, only: continuity{0}_code".format(tag)
in code)
assert "call continuity{0}_code(".format(tag) in code
# The kernel and module name should have gained the tag just identified
# and be written to the CWD
filename = os.path.join(str(tmpdir), "continuity{0}_mod.f90".format(tag))
assert os.path.isfile(filename)
# Parse the new kernel file
f2003_parser = parser.ParserFactory().create()
reader = FortranFileReader(filename)
prog = f2003_parser(reader)
# Check that the module has the right name
modules = walk_ast(prog.content, [Fortran2003.Module_Stmt])
assert str(modules[0].items[1]) == "continuity{0}_mod".format(tag)
# Check that the subroutine has the right name
subs = walk_ast(prog.content, [Fortran2003.Subroutine_Stmt])
found = False
for sub in subs:
if str(sub.items[1]) == "continuity{0}_code".format(tag):
found = True
break
assert found
# Check that the kernel type has been re-named
dtypes = walk_ast(prog.content, [Fortran2003.Derived_Type_Def])
names = walk_ast(dtypes[0].content, [Fortran2003.Type_Name])
assert str(names[0]) == "continuity{0}_type".format(tag)
from gocean1p0_build import GOcean1p0Build
# If compilation fails this will raise an exception
GOcean1p0Build(tmpdir).compile_file(filename)
old_cwd.chdir()
def test_use_stmts():
''' Test that generating code for OpenCL results in the correct
module use statements. '''
psy, _ = get_invoke("single_invoke.f90", API, idx=0)
sched = psy.invokes.invoke_list[0].schedule
otrans = OCLTrans()
otrans.apply(sched)
generated_code = str(psy.gen).lower()
expected = '''\
subroutine invoke_0_compute_cu(cu_fld, p_fld, u_fld)
use fortcl, only: create_rw_buffer
use fortcl, only: get_num_cmd_queues, get_cmd_queues, get_kernel_by_name
use clfortran
use iso_c_binding'''
assert expected in generated_code
assert "if (first_time) then" in generated_code
def test_new_kernel_dir(tmpdir, monkeypatch):
''' Check that we write out the transformed kernel to a specified
directory. '''
# Set the output directory in the configuration object
config = Config.get()
monkeypatch.setattr(config, "_kernel_output_dir", str(tmpdir))
psy, invoke = get_invoke("nemolite2d_alg_mod.f90", api="gocean1.0", idx=0)
sched = invoke.schedule
kern = sched.children[0].children[0].children[0]
rtrans = ACCRoutineTrans()
_, _ = rtrans.apply(kern)
# Generate the code (this triggers the generation of a new kernel)
_ = str(psy.gen)
file_list = os.listdir(str(tmpdir))
assert len(file_list) == 1
assert file_list[0] == 'continuity_0_mod.f90'
def test_goloop():
''' Check the handling of non-NEMO do loops.
TODO #440: Does not work atm, GOLoops also have start/stop as
strings, which are even not defined. Only after genCode is called will
they be defined.
'''
_, invoke = get_invoke("single_invoke_two_kernels_scalars.f90",
"gocean1.0",
name="invoke_0")
do_loop = invoke.schedule.children[0]
assert isinstance(do_loop, Loop)
var_accesses = VariablesAccessInfo()
do_loop.reference_accesses(var_accesses)
assert str(var_accesses) == ": READ, a_scalar: READ, i: READ+WRITE, "\
"j: READ+WRITE, " "ssh_fld: READ+WRITE, "\
"tmask: READ"
def test_profile_kernels_gocean1p0():
'''Check that all kernels are instrumented correctly
'''
Profiler.set_options([Profiler.KERNELS])
_, invoke = get_invoke("single_invoke_two_kernels.f90", "gocean1.0",
idx=0)
# Convert the invoke to code, and remove all new lines, to make
# regex matching easier
code = str(invoke.gen()).replace("\n", "")
# Test that kernel profiling works in case of two kernel calls
# in a single invoke subroutine - i.e. we need to have one profile
# start call before two nested loops, and one profile end call
# after that.
# Also note that the '.*' after compute_cu_code is necessary since
# the name could be changed to avoid duplicates (depending on order
# in which the tests are executed).
correct_re = ("subroutine invoke.*"
"use profile_mod, only: ProfileData.*"
r"TYPE\(ProfileData\), save :: profile.*"
r"TYPE\(ProfileData\), save :: profile.*"
r"call ProfileStart\(\"compute_cu_mod\", "
r"\"compute_cu_code.*\", (?P<profile1>\w*)\).*"
"do j.*"
"do i.*"
"call.*"
"end.*"
"end.*"
r"call ProfileEnd\((?P=profile1)\).*"
r"call ProfileStart\(\"time_smooth_mod\", "
r"\"time_smooth_code\", (?P<profile2>\w*)\).*"
"do j.*"
"do i.*"
"call.*"
"end.*"
"end.*"
r"call ProfileEnd\((?P=profile2)\)")
groups = re.search(correct_re, code, re.I)
assert groups is not None
# Check that the variables are different
assert groups.group(1) != groups.group(2)
Profiler.set_options(None)
def test_kern_case_insensitive(mod_name, sub_name, kernel_outputdir,
monkeypatch):
'''Check that the test to see if a kernel conforms to the <name>_mod,
<name>_code convention is case insensitive. This check also tests that the
removal of _mod to create part of the output filename is case
insensitive.
'''
_, invoke = get_invoke("1_single_invoke.f90", api="dynamo0.3", idx=0)
sched = invoke.schedule
kernels = sched.walk(Kern)
kern = kernels[0]
ktrans = Dynamo0p3KernelConstTrans()
_, _ = ktrans.apply(kern, number_of_layers=100)
monkeypatch.setattr(kern, "_module_name", mod_name)
monkeypatch.setattr(kern, "_name", sub_name)
# Generate the code - this should not raise an exception.
kern.rename_and_write()
filename = os.path.join(str(kernel_outputdir), mod_name[:8] + "_0_mod.f90")
assert os.path.isfile(filename)
def test_goloop_partially():
''' Check the handling of non-NEMO do loops.
TODO #440: This test is identical to test_goloop above, but it asserts in a
way that works before #440 is fixed, so that we make sure we test the rest
of the gocean variable access handling.
'''
_, invoke = get_invoke("single_invoke_two_kernels_scalars.f90",
"gocean1.0",
name="invoke_0")
do_loop = invoke.schedule.children[0]
assert isinstance(do_loop, Loop)
# The third argument is GO_GRID_X_MAX_INDEX, which is scalar
assert do_loop.args[2].is_scalar()
# The fourth argument is GO_GRID_MASK_T, which is an array
assert not do_loop.args[3].is_scalar()
var_accesses = VariablesAccessInfo()
do_loop.reference_accesses(var_accesses)
assert "a_scalar: READ, i: READ+WRITE, j: READ+WRITE, "\
"ssh_fld: READWRITE, ssh_fld%grid%subdomain%internal%xstop: READ, "\
"ssh_fld%grid%tmask: READ" in str(var_accesses)
