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