def test_incorrect_num_of_args():
src = "kernel void A(const int a) {}"
# too many inputs
with test.Raises(ValueError):
driver.DriveKernel(
env.OclgrindOpenCLEnvironment(),
src,
[[1], [2], [3]],
gsize=(1, 1, 1),
lsize=(1, 1, 1),
)
# too few inputs
with test.Raises(ValueError):
driver.DriveKernel(env.OclgrindOpenCLEnvironment(),
src, [],
gsize=(1, 1, 1),
lsize=(1, 1, 1))
# incorrect input width (3 ints instead of one)
with test.Raises(ValueError):
driver.DriveKernel(
env.OclgrindOpenCLEnvironment(),
src,
[[1, 2, 3]],
gsize=(1, 1, 1),
lsize=(1, 1, 1),
)
def test_ClLauncherHarness_oclgrind_testbed():
"""Test that harness can be made from project-local oclgrind."""
config = harness_pb2.ClLauncherHarness()
config.opencl_env.extend([
env.OclgrindOpenCLEnvironment().name,
env.OclgrindOpenCLEnvironment().name
])
config.opencl_opt.extend([True, False])
harness = cl_launcher.ClLauncherHarness(config)
assert len(harness.testbeds) == 2
assert harness.testbeds[0].name == env.OclgrindOpenCLEnvironment().name
assert harness.testbeds[0].opts["opencl_opt"] == "enabled"
assert harness.testbeds[1].name == env.OclgrindOpenCLEnvironment().name
assert harness.testbeds[1].opts["opencl_opt"] == "disabled"
def test_empty_kernel():
src = " kernel void A() {} "
outputs = driver.DriveKernel(env.OclgrindOpenCLEnvironment(),
src, [],
gsize=(1, 1, 1),
lsize=(1, 1, 1))
assert len(outputs) == 0
def test_rewrite_compile_link_execute_clinfo(tempdir: pathlib.Path,
clinfo_src: str):
log = _RewriteCompileLinkExecute(
tempdir,
clinfo_src,
lang="c",
extra_ldflags=["-lm", "-lstdc++"],
extra_exec_args=["--raw"],
)
print(log)
assert log.ms_since_unix_epoch
assert log.returncode == 0
assert log.device == cldrive_env.OclgrindOpenCLEnvironment().proto
assert len(log.kernel_invocation) == 0
assert len(log.opencl_program_source) == 0
assert not log.stderr
assert re.match(
r"0 CL_PLATFORM_NAME Oclgrind\n"
r"0 CL_PLATFORM_VERSION OpenCL \d+\.\d+ \(Oclgrind [\d\.]+\)\n"
r"0:0 CL_DEVICE_NAME Oclgrind Simulator\n"
r"0:0 CL_DEVICE_TYPE [a-zA-Z |]+\n"
r"0:0 CL_DEVICE_VERSION OpenCL \d+\.\d+ \(Oclgrind [\d\.]+\)\n"
r"0:0 CL_DEVICE_GLOBAL_MEM_SIZE \d+\n"
r"0:0 CL_DEVICE_LOCAL_MEM_SIZE \d+\n"
r"0:0 CL_DEVICE_MAX_WORK_GROUP_SIZE \d+\n"
r"0:0 CL_DEVICE_MAX_WORK_ITEM_SIZES \(\d+, \d+, \d+\)\n",
log.stdout,
re.MULTILINE,
)
def GetGoldStandardTestHarness() -> base_harness.HarnessBase:
"""Instantiate the gold standard test harness.
Uses the global FLAGS to determine the harness to instantiate.
Returns:
A Harness instance.
"""
if FLAGS.generator == "clgen":
harness_class = cldrive.CldriveHarness
config_class = harness_pb2.CldriveHarness
elif FLAGS.generator == "clsmith":
harness_class = cl_launcher.ClLauncherHarness
config_class = harness_pb2.ClLauncherHarness
else:
raise app.UsageError(
f"Unrecognized value for --generator: '{FLAGS.generator}'"
)
app.Log(1, "Preparing gold standard testbed.")
config = GetBaseHarnessConfig(config_class)
config.opencl_env.extend([env.OclgrindOpenCLEnvironment().name])
config.opencl_opt.extend([True])
gs_harness = harness_class(config)
assert len(gs_harness.testbeds) >= 1
return gs_harness
def test_gsize_smaller_than_lsize():
src = "kernel void A() {}"
with test.Raises(ValueError):
driver.DriveKernel(env.OclgrindOpenCLEnvironment(),
src, [],
gsize=(4, 1, 1),
lsize=(8, 1, 1))
def test_RewriteClDeviceType_rewrites_file(tempdir: pathlib.Path):
"""Test that CL_DEVICE_TYPE is rewritten in file."""
with open(tempdir / "foo", "w") as f:
f.write("Hello world! The device type is: CL_DEVICE_TYPE_GPU.")
gpgpu.RewriteClDeviceType(cldrive_env.OclgrindOpenCLEnvironment(), tempdir)
with open(tempdir / "foo") as f:
assert f.read() == "Hello world! The device type is: CL_DEVICE_TYPE_CPU."
def test_CldriveHarness_oclgrind_testbed_count_two():
"""Test that correct number of testbeds are instantiated."""
oclgrind_env_name = env.OclgrindOpenCLEnvironment().name
config = harness_pb2.CldriveHarness()
config.opencl_env.extend([oclgrind_env_name, oclgrind_env_name])
config.opencl_opt.extend([True, False])
harness = cldrive.CldriveHarness(config)
assert len(harness.testbeds) == 2
def test_zero_size_input():
src = "kernel void A(global int* a) {}"
with test.Raises(ValueError):
driver.DriveKernel(
env.OclgrindOpenCLEnvironment(),
src,
[[]],
gsize=(1, 1, 1),
lsize=(1, 1, 1),
)
def test_syntax_error():
src = "kernel void A(gl ob a l i nt* a) {}"
with testlib.DevNullRedirect():
with test.Raises(driver.OpenCLValueError):
driver.DriveKernel(
env.OclgrindOpenCLEnvironment(),
src,
[[]],
gsize=(1, 1, 1),
lsize=(1, 1, 1),
)
def test_CldriveHarness_oclgrind_testbed_names():
"""Test that correct names set on testbeds."""
oclgrind_env_name = env.OclgrindOpenCLEnvironment().name
config = harness_pb2.CldriveHarness()
config.opencl_env.extend([oclgrind_env_name, oclgrind_env_name])
config.opencl_opt.extend([True, False])
harness = cldrive.CldriveHarness(config)
assert harness.testbeds[0].name == oclgrind_env_name
assert harness.testbeds[1].name == oclgrind_env_name
def test_CldriveHarness_oclgrind_testbed_opts():
"""Test that opencl_opt option set on testbeds."""
oclgrind_env_name = env.OclgrindOpenCLEnvironment().name
config = harness_pb2.CldriveHarness()
config.opencl_env.extend([oclgrind_env_name, oclgrind_env_name])
config.opencl_opt.extend([True, False])
harness = cldrive.CldriveHarness(config)
assert harness.testbeds[0].opts["opencl_opt"] == "enabled"
assert harness.testbeds[1].opts["opencl_opt"] == "disabled"
def test_timeout():
# non-terminating kernel
src = "kernel void A() { while (true) ; }"
with test.Raises(driver.Timeout):
driver.DriveKernel(
env.OclgrindOpenCLEnvironment(),
src,
[],
gsize=(1, 1, 1),
lsize=(1, 1, 1),
timeout=1,
)
def test_OclgrindOpenCLEnvironment_Exec_version():
"""Test that OclgrindOpenCLEnvironment.Exec() works as expected."""
proc = env.OclgrindOpenCLEnvironment().Exec(["--version"])
# This test will of course fail if the @oclgrind package is updated.
assert (proc.stdout == """
Oclgrind 18.3
Copyright (c) 2013-2018
James Price and Simon McIntosh-Smith, University of Bristol
https://github.com/jrprice/Oclgrind
""")
def test_data_unchanged():
src = "kernel void A(global int* a, global int* b, const int c) {}"
inputs = data.MakeRand(src, 16)
outputs = driver.DriveKernel(
env.OclgrindOpenCLEnvironment(),
src,
inputs,
gsize=(16, 1, 1),
lsize=(1, 1, 1),
)
testlib.Assert2DArraysAlmostEqual(outputs, inputs)
def test_invalid_sizes():
src = "kernel void A() {}"
# invalid global size
with test.Raises(ValueError):
driver.DriveKernel(
env.OclgrindOpenCLEnvironment(),
src,
[],
gsize=(0, -4, 1),
lsize=(1, 1, 1),
)
# invalid local size
with test.Raises(ValueError):
driver.DriveKernel(
env.OclgrindOpenCLEnvironment(),
src,
[],
gsize=(1, 1, 1),
lsize=(-1, 1, 1),
)
def test_vector_input():
inputs = [[0, 1, 2, 3, 0, 1, 2, 3], [2, 4]]
inputs_orig = [[0, 1, 2, 3, 0, 1, 2, 3], [2, 4]]
outputs_gs = [[0, 2, 4, 6, 0, 4, 8, 12], [2, 4]]
src = """
kernel void A(global int* a, const int2 b) {
const int x_id = get_global_id(0);
const int y_id = get_global_id(1);
if (!y_id) {
a[x_id] *= b.x;
} else {
a[get_global_size(0) + x_id] *= b.y;
}
}
"""
outputs = driver.DriveKernel(
env.OclgrindOpenCLEnvironment(),
src,
inputs,
gsize=(4, 2, 1),
lsize=(1, 1, 1),
)
testlib.Assert2DArraysAlmostEqual(inputs, inputs_orig)
testlib.Assert2DArraysAlmostEqual(outputs, outputs_gs)
# run kernel a second time with the previous outputs
outputs2 = driver.DriveKernel(
env.OclgrindOpenCLEnvironment(),
src,
outputs,
gsize=(4, 2, 1),
lsize=(1, 1, 1),
)
outputs2_gs = [[0, 4, 8, 12, 0, 16, 32, 48], [2, 4]]
testlib.Assert2DArraysAlmostEqual(outputs2, outputs2_gs)
def test_CldriveHarness_oclgrind_testbed_uneven_name_and_opt():
"""Error is raised if number of opt_opt != number of opencl_env."""
oclgrind_env_name = env.OclgrindOpenCLEnvironment().name
config = harness_pb2.CldriveHarness()
config.opencl_env.extend([oclgrind_env_name, oclgrind_env_name])
config.opencl_opt.extend([True])
with test.Raises(ValueError) as e_ctx:
cldrive.CldriveHarness(config, default_to_all_environments=False)
assert (
"CldriveHarness.opencl_env and CldriveHarness.opencl_opt lists are "
"not the same length") in str(e_ctx.value)
def test_iterative_increment():
src = "kernel void A(global int* a) { a[get_global_id(0)] += 1; }"
d_cl, d_host = [np.arange(16)], np.arange(16)
for _ in range(8):
d_host += 1 # perform computation on host
d_cl = driver.DriveKernel(
env.OclgrindOpenCLEnvironment(),
src,
d_cl,
gsize=(16, 1, 1),
lsize=(16, 1, 1),
)
testlib.Assert2DArraysAlmostEqual(d_cl, [d_host])
def test_data_zerod():
# zero-ing a randomly initialized array
src = "kernel void A(global int* a) { a[get_global_id(0)] = 0; }"
inputs = data.MakeRand(src, 16)
outputs = driver.DriveKernel(
env.OclgrindOpenCLEnvironment(),
src,
inputs,
gsize=(16, 1, 1),
lsize=(4, 1, 1),
)
testlib.Assert2DArraysAlmostEqual(outputs, [np.zeros(16)])
def test_ExecClsmithSource_syntax_error():
"""Test outcome of kernel with syntax error."""
env_ = env.OclgrindOpenCLEnvironment()
proc = cl_launcher.ExecClsmithSource(
env_,
"!@!###syntax error!",
driver.NDRange(1, 1, 1),
driver.NDRange(1, 1, 1),
"---debug",
)
assert proc.returncode == 1
assert proc.stdout == ""
assert "Error building program: -11" in proc.stderr
def test_gsize_smaller_than_data():
src = "kernel void A(global int* a) { a[get_global_id(0)] = 0; }"
inputs = [[5, 5, 5, 5, 5, 5, 5, 5]]
outputs_gs = [[0, 0, 0, 0, 5, 5, 5, 5]]
outputs = driver.DriveKernel(
env.OclgrindOpenCLEnvironment(),
src,
inputs,
gsize=(4, 1, 1),
lsize=(4, 1, 1),
)
testlib.Assert2DArraysAlmostEqual(outputs, outputs_gs)
def test_comparison_against_pointer_warning():
src = """
kernel void A(global int* a) {
int id = get_global_id(0);
if (id < a) a += 1;
}
"""
driver.DriveKernel(
env.OclgrindOpenCLEnvironment(),
src,
[[0]],
gsize=(1, 1, 1),
lsize=(1, 1, 1),
)
def test_BenchmarkSuite_integration_test(benchmark_suite: typing.Callable,
tempdir: pathlib.Path):
"""Test compilation and execution of benchmark suite using oclgrind."""
with benchmark_suite() as bs:
bs.ForceOpenCLEnvironment(cldrive_env.OclgrindOpenCLEnvironment())
observer = MockBenchmarkObserver(stop_after=1)
# `stop_after` raises BenchmarkInterrupt.
try:
bs.Run([observer])
assert False
except gpgpu.BenchmarkInterrupt:
pass
assert len(observer.logs) == 1
assert observer.logs[0].benchmark_name in bs.benchmarks
def test_ExecClsmithSource_pass():
"""And end-to-end test of executing a CLSmith source."""
env_ = env.OclgrindOpenCLEnvironment()
proc = cl_launcher.ExecClsmithSource(
env_,
CLSMITH_EXAMPLE_SRC,
driver.NDRange(1, 1, 1),
driver.NDRange(1, 1, 1),
"---debug",
)
assert not proc.returncode
assert "3-D global size 1 = [1, 1, 1]" in proc.stderr
assert "3-D local size 1 = [1, 1, 1]" in proc.stderr
assert "OpenCL optimizations: on" in proc.stderr
assert "Platform: " in proc.stderr
assert "Device: " in proc.stderr
assert "Compilation terminated successfully..."
assert proc.stdout == "0,"
def _RewriteCompileLinkExecute(
outdir: pathlib.Path,
src: str,
lang: str = "c++",
extra_ldflags=None,
extra_cflags=None,
extra_exec_args=None,
) -> libcecl_pb2.LibceclExecutableRun:
"""Compile, link, and execute a program using libcecl."""
# Re-write OpenCL source to use libcecl.
libcecl_src = libcecl_rewriter.RewriteOpenClSource(src)
# Compile libcecl source to bytecode.
src_path = outdir / f"a.txt"
objectfile_path = outdir / "a.o"
cflags, ldflags = libcecl_compile.LibCeclCompileAndLinkFlags()
with open(src_path, "w") as f:
f.write(libcecl_src)
extra_cflags = extra_cflags or []
subprocess.check_call(
["c++", "-x", lang,
str(src_path), "-c", "-o",
str(objectfile_path)] + cflags + extra_cflags)
assert objectfile_path.is_file()
# Compile bytecode to executable and link.
bin_path = outdir / "a.out"
extra_ldflags = extra_ldflags or []
subprocess.check_call(
["c++", "-o", str(bin_path),
str(objectfile_path)] + ldflags + extra_ldflags)
assert bin_path.is_file()
# Run executable on oclgrind.
extra_exec_args = extra_exec_args or []
return libcecl_runtime.RunLibceclExecutable(
[str(bin_path)] + extra_exec_args,
cldrive_env.OclgrindOpenCLEnvironment())
def test_rewrite_compile_link_execute(tempdir: pathlib.Path, hello_src: str):
"""Test end-to-end libcecl pipeline."""
log = _RewriteCompileLinkExecute(tempdir,
hello_src,
lang="c++",
extra_cflags=["-std=c++11"])
print(log)
# Check values in log.
assert log.ms_since_unix_epoch
assert log.returncode == 0
assert log.device == cldrive_env.OclgrindOpenCLEnvironment().proto
assert len(log.kernel_invocation) == 1
assert len(log.opencl_program_source) == 1
assert (log.opencl_program_source[0] == """\
__kernel void vadd(
__global float* a,
__global float* b,
__global float* c,
const unsigned int count)
{
int i = get_global_id(0);
if(i < count)
c[i] = a[i] + b[i];
}""")
assert log.kernel_invocation[0].kernel_name == "vadd"
assert log.kernel_invocation[0].global_size == 1024
assert log.kernel_invocation[0].local_size == 0
assert log.kernel_invocation[0].transferred_bytes == 12288
assert (log.kernel_invocation[0].transfer_time_ns > 1000
) # Flaky, but probably true.
assert (log.kernel_invocation[0].kernel_time_ns > 1000
) # Flaky, but probably true.
profile_time = (log.kernel_invocation[0].transfer_time_ns +
log.kernel_invocation[0].kernel_time_ns)
assert 1000 < profile_time < log.elapsed_time_ns
def test_profiling():
src = """
kernel void A(global int* a, constant int* b) {
const int id = get_global_id(0);
a[id] *= b[id];
}
"""
inputs = [np.arange(16), np.arange(16)]
outputs_gs = [np.arange(16)**2, np.arange(16)]
with testlib.DevNullRedirect():
outputs = driver.DriveKernel(
env.OclgrindOpenCLEnvironment(),
src,
inputs,
gsize=(16, 1, 1),
lsize=(16, 1, 1),
profiling=True,
)
testlib.Assert2DArraysAlmostEqual(outputs, outputs_gs)
def test_simple():
inputs = [[0, 1, 2, 3, 4, 5, 6, 7]]
inputs_orig = [[0, 1, 2, 3, 4, 5, 6, 7]]
outputs_gs = [[0, 2, 4, 6, 8, 10, 12, 14]]
src = """
kernel void A(global float* a) {
const int x_id = get_global_id(0);
a[x_id] *= 2.0;
}
"""
outputs = driver.DriveKernel(
env.OclgrindOpenCLEnvironment(),
src,
inputs,
gsize=(8, 1, 1),
lsize=(1, 1, 1),
)
testlib.Assert2DArraysAlmostEqual(inputs, inputs_orig)
testlib.Assert2DArraysAlmostEqual(outputs, outputs_gs)
def device() -> env.OpenCLEnvironment: """Test fixture which yields a testing OpenCL device.""" return env.OclgrindOpenCLEnvironment().proto
