def test_oclw_kernelschedule():
'''Check the OpenCLWriter class kernelschedule_node visitor produces
the expected C code.
'''
# The kernelschedule OpenCL Backend relies on abstrct methods that
# need to be implemented by the APIs. A generic kernelschedule will
# produce a NotImplementedError.
oclwriter = OpenCLWriter()
kschedule = KernelSchedule("kname")
with pytest.raises(NotImplementedError) as error:
_ = oclwriter(kschedule)
assert "Abstract property. Which symbols are data arguments is " \
"API-specific." in str(error)
# Mock abstract properties. (pytest monkeypatch does not work
# with properties, used sub-class instead)
class MockSymbolTable(SymbolTable):
''' Mock needed abstract methods of the Symbol Table '''
@property
def iteration_indices(self):
return self.argument_list[:2]
@property
def data_arguments(self):
return self.argument_list[2:]
kschedule.symbol_table.__class__ = MockSymbolTable
# Create a sample symbol table and kernel schedule
interface = Symbol.Argument(access=Symbol.Access.UNKNOWN)
i = Symbol('i', 'integer', interface=interface)
j = Symbol('j', 'integer', interface=interface)
data1 = Symbol('data1', 'real', [10, 10], interface=interface)
data2 = Symbol('data2', 'real', [10, 10], interface=interface)
kschedule.symbol_table.add(i)
kschedule.symbol_table.add(j)
kschedule.symbol_table.add(data1)
kschedule.symbol_table.add(data2)
kschedule.symbol_table.specify_argument_list([i, j, data1, data2])
kschedule.addchild(Return(parent=kschedule))
result = oclwriter(kschedule)
print(result)
assert result == "" \
"__kernel void kname(\n" \
" __global double * restrict data1,\n" \
" __global double * restrict data2\n" \
" ){\n" \
" int data1LEN1 = get_global_size(0);\n" \
" int data1LEN2 = get_global_size(1);\n" \
" int data2LEN1 = get_global_size(0);\n" \
" int data2LEN2 = get_global_size(1);\n" \
" int i = get_global_id(0);\n" \
" int j = get_global_id(1);\n" \
" return;\n" \
"}\n"
def test_oclw_initialization():
'''Test that the OpenCLWriter-specific parameters are error checked'''
# OCLWriter can be initialized with default values
oclwriter = OpenCLWriter()
assert oclwriter._kernels_local_size == 1
# Pass a kernels_local_size parameter
with pytest.raises(TypeError) as error:
oclwriter = OpenCLWriter(kernels_local_size='invalid')
assert "kernel_local_size should be an integer but found 'str'." \
in str(error)
with pytest.raises(ValueError) as error:
oclwriter = OpenCLWriter(kernels_local_size=-4)
assert "kernel_local_size should be a positive integer but found -4." \
in str(error)
oclwriter = OpenCLWriter(kernels_local_size=4)
assert oclwriter._kernels_local_size == 4
def test_oclw_gen_declaration():
'''Check the OpenCLWriter class gen_declaration method produces
the expected declarations.
'''
oclwriter = OpenCLWriter()
# Basic entry - Scalar are passed by value and don't have additional
# qualifiers.
symbol = Symbol("dummy1", "integer")
result = oclwriter.gen_declaration(symbol)
assert result == "int dummy1"
# Array argument has a memory qualifier (only __global for now)
symbol = Symbol("dummy2", "integer", shape=[2, None, 2])
result = oclwriter.gen_declaration(symbol)
assert result == "__global int * restrict dummy2"
# Array with unknown intent
symbol = Symbol("dummy2", "integer", shape=[2, None, 2],
interface=Symbol.Argument(access=Symbol.Access.UNKNOWN))
result = oclwriter.gen_declaration(symbol)
assert result == "__global int * restrict dummy2"
def test_opencl_code_generation_with_boundary_mask():
''' Tests that OpenCL kernel generated after applying the
GOMoveIterationBoundariesInsideKernelTrans has the 4 boundary values as
kernel arguments and has a masking statement at the beginning of the
executable code.
'''
psy, _ = get_invoke("single_invoke.f90", API, idx=0, dist_mem=False)
sched = psy.invokes.invoke_list[0].schedule
kernel = sched.children[0].loop_body[0].loop_body[0] # compute_cu kernel
trans = GOMoveIterationBoundariesInsideKernelTrans()
trans.apply(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"
" int xstart,\n"
" int xstop,\n"
" int ystart,\n"
" int ystop\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"
" if ((((i < xstart) || (i > xstop)) || ((j < ystart) ||"
" (j > ystop)))) {\n"
" return;\n"
" }\n"
" cu[j * cuLEN1 + i] = ((0.5e0 * (p[j * pLEN1 + (i + 1)]"
" + p[j * pLEN1 + i])) * u[j * uLEN1 + i]);\n"
"}\n\n")
openclwriter = OpenCLWriter()
assert expected_code == openclwriter(kschedule)
def test_oclw_gen_array_length_variables():
'''Check the OpenCLWriter class gen_array_length_variables method produces
the expected declarations.
'''
oclwriter = OpenCLWriter()
# A scalar should not return any LEN variables
symbol1 = DataSymbol("dummy2LEN1", INTEGER_TYPE)
result = oclwriter.gen_array_length_variables(symbol1)
assert result == ""
# Array with 1 dimension generates 1 length variable
array_type = ArrayType(INTEGER_TYPE, [2])
symbol2 = DataSymbol("dummy1", array_type)
result = oclwriter.gen_array_length_variables(symbol2)
assert result == "int dummy1LEN1 = get_global_size(0);\n"
# Array with multiple dimension generates one variable per dimension
array_type = ArrayType(INTEGER_TYPE, [2, ArrayType.Extent.ATTRIBUTE, 2])
symbol3 = DataSymbol("dummy2", array_type)
result = oclwriter.gen_array_length_variables(symbol3)
assert result == "int dummy2LEN1 = get_global_size(0);\n" \
"int dummy2LEN2 = get_global_size(1);\n" \
"int dummy2LEN3 = get_global_size(2);\n"
# Create a symbol table
symtab = SymbolTable()
symtab.add(symbol1)
symtab.add(symbol2)
symtab.add(symbol3)
# If there are no name clashes, generate array length variables.
result = oclwriter.gen_array_length_variables(symbol2, symtab)
assert result == "int dummy1LEN1 = get_global_size(0);\n"
with pytest.raises(VisitorError) as excinfo:
_ = oclwriter.gen_array_length_variables(symbol3, symtab)
assert "Unable to declare the variable 'dummy2LEN1' to store the length " \
"of 'dummy2' because the Symbol Table already contains a symbol with" \
" the same name." in str(excinfo.value)
def test_oclw_kernelschedule():
'''Check the OpenCLWriter class kernelschedule_node visitor produces
the expected OpenCL code.
'''
# The kernelschedule OpenCL Backend relies on abstract methods that
# need to be implemented by the APIs. A generic kernelschedule will
# produce a NotImplementedError.
oclwriter = OpenCLWriter()
kschedule = KernelSchedule("kname")
with pytest.raises(NotImplementedError) as error:
_ = oclwriter(kschedule)
assert "Abstract property. Which symbols are data arguments is " \
"API-specific." in str(error.value)
# Mock abstract properties. (pytest monkeypatch does not work
# with properties, used sub-class instead)
class MockSymbolTable(SymbolTable):
''' Mock needed abstract methods of the Symbol Table '''
@property
def iteration_indices(self):
return self.argument_list[:2]
@property
def data_arguments(self):
return self.argument_list[2:]
kschedule.symbol_table.__class__ = MockSymbolTable
# Create a sample symbol table and kernel schedule
interface = ArgumentInterface(ArgumentInterface.Access.UNKNOWN)
i = DataSymbol('i', INTEGER_TYPE, interface=interface)
j = DataSymbol('j', INTEGER_TYPE, interface=interface)
array_type = ArrayType(REAL_TYPE, [10, 10])
data1 = DataSymbol('data1', array_type, interface=interface)
data2 = DataSymbol('data2', array_type, interface=interface)
kschedule.symbol_table.add(i)
kschedule.symbol_table.add(j)
kschedule.symbol_table.add(data1)
kschedule.symbol_table.add(data2)
kschedule.symbol_table.specify_argument_list([i, j, data1, data2])
kschedule.addchild(Return(parent=kschedule))
result = oclwriter(kschedule)
assert result == "" \
"__kernel void kname(\n" \
" __global double * restrict data1,\n" \
" __global double * restrict data2\n" \
" ){\n" \
" int data1LEN1 = get_global_size(0);\n" \
" int data1LEN2 = get_global_size(1);\n" \
" int data2LEN1 = get_global_size(0);\n" \
" int data2LEN2 = get_global_size(1);\n" \
" int i = get_global_id(0);\n" \
" int j = get_global_id(1);\n" \
" return;\n" \
"}\n\n"
# Set a local_size value different to 1 into the KernelSchedule
oclwriter = OpenCLWriter(kernels_local_size=4)
result = oclwriter(kschedule)
assert result == "" \
"__attribute__((reqd_work_group_size(4, 1, 1)))\n" \
"__kernel void kname(\n" \
" __global double * restrict data1,\n" \
" __global double * restrict data2\n" \
" ){\n" \
" int data1LEN1 = get_global_size(0);\n" \
" int data1LEN2 = get_global_size(1);\n" \
" int data2LEN1 = get_global_size(0);\n" \
" int data2LEN2 = get_global_size(1);\n" \
" int i = get_global_id(0);\n" \
" int j = get_global_id(1);\n" \
" return;\n" \
"}\n\n"
# Add a symbol with a deferred interface and check that this raises the
# expected error
array_type = ArrayType(REAL_TYPE, [10, 10])
kschedule.symbol_table.add(
DataSymbol('broken', array_type, interface=UnresolvedInterface()))
with pytest.raises(VisitorError) as err:
_ = oclwriter(kschedule)
assert ("symbol table contains unresolved data entries (i.e. that have no "
"defined Interface) which are not used purely to define the "
"precision of other symbols: 'broken'" in str(err.value))