def test_GetKernelArguments_multiple_kernels():
"""Test that error is raised if no kernel defined."""
with pytest.raises(args.MultipleKernelsError):
args.GetKernelArguments("""
kernel void A() {}
kernel void B() {}
""")
def test_GetKernelArguments_properties():
"""Test extracted properties of kernel."""
args_ = args.GetKernelArguments("""
kernel void A(const global int* a, global const float* b,
local float4 *const c, const int d, float2 e) {}
""")
assert len(args_) == 5
assert args_[0].is_pointer
assert args_[0].address_space == "global"
assert args_[0].typename == "int"
assert args_[0].name == "a"
assert args_[0].bare_type == "int"
assert not args_[0].is_vector
assert args_[0].vector_width == 1
assert args_[0].is_const
assert args_[1].is_pointer
assert args_[1].address_space == "global"
assert args_[1].typename == "float"
assert args_[1].name == "b"
assert args_[1].bare_type == "float"
assert not args_[1].is_vector
assert args_[1].vector_width == 1
assert args_[1].is_const
assert args_[2].is_pointer
assert args_[2].address_space == "local"
assert args_[2].typename == "float4"
assert args_[2].name == "c"
assert args_[2].bare_type == "float"
assert args_[2].is_vector
assert args_[2].vector_width == 4
assert not args_[2].is_const
def test_GetKernelArguments_address_spaces():
"""Test address space types."""
args_ = args.GetKernelArguments("""
kernel void A(global int* a,
local int* b,
constant int* c,
const int d) {}
""")
assert len(args_) == 4
assert args_[0].address_space == "global"
assert args_[1].address_space == "local"
assert args_[2].address_space == "constant"
assert args_[3].address_space == "private"
def test_GetKernelArguments_proper_kernel():
"""Test kernel arguments of a "real" kernel."""
args_ = args.GetKernelArguments("""
typedef int foobar;
void B(const int e);
__kernel void A(const __global int* data, __local float4 * restrict car,
__global const float* b, const int foo, int d) {
int tid = get_global_id(0);
data[tid] *= 2.0;
}
void B(const int e) {}
""")
assert len(args_) == 5
assert args_[0].is_const
assert args_[0].is_pointer
assert args_[0].typename == "int"
assert args_[0].bare_type == "int"
def emit_c(src: str, inputs: np.array, gsize: typing.Optional[driver.NDRange],
lsize: typing.Optional[driver.NDRange], timeout: int = -1,
optimizations: bool = True, profiling: bool = False,
debug: bool = False, compile_only: bool = False,
create_kernel: bool = True) -> np.array:
"""
Generate C code to drive kernel.
Parameters
----------
env : OpenCLEnvironment
The OpenCL environment to run the kernel in.
src : str
The OpenCL kernel source.
inputs : np.array
The input data to the kernel.
optimizations : bool, optional
Whether to enable or disbale OpenCL compiler optimizations.
profiling : bool, optional
If true, print OpenCLevent times for data transfers and kernel
executions to stderr.
timeout : int, optional
Cancel execution if it has not completed after this many seconds.
A value <= 0 means never time out.
debug : bool, optional
If true, silence the OpenCL compiler.
compile_only: bool, optional
If true, generate code only to compile the kernel, not to generate
inputs and run it.
create_kernel: bool, optional
If 'compile_only' parameter is set, this parameter determines whether
to create a kernel object after compilation. This requires a kernel
name.
Returns
-------
str
Code which can be compiled using a C compiler to drive the kernel.
Raises
------
ValueError
If input types are incorrect.
TypeError
If an input is of an incorrect type.
LogicError
If the input types do not match OpenCL kernel types.
PorcelainError
If the OpenCL subprocess exits with non-zero return code.
RuntimeError
If OpenCL program fails to build or run.
Examples
--------
TODO
"""
src_string = escape_c_string(src)
optimizations_on_off = "on" if optimizations else "off"
clBuildProgram_opts = "NULL" if optimizations else '"-cl-opt-disable"'
c = f"""
/*
* Usage:
* gcc -std=c99 [-DPLATFORM_ID=<platform-id>] [-DDEVICE_ID=<device-id>] foo.c -lOpenCL
* ./a.out [-p <platform-id>] [-d <device-id>]
*
* Host code generated using cldrive <https://github.com/ChrisCummins/cldrive>
*/
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif
const char *kernel_src = \\
{src_string};
#ifndef PLATFORM_ID
# define PLATFORM_ID 0
#endif
#ifndef DEVICE_ID
# define DEVICE_ID 0
#endif
#define True 1
#define False 0
#ifndef __APPLE__
typedef unsigned char bool;
#endif
typedef unsigned short ushort;
const char *clerror_string(cl_int err) {{
/* written by @Selmar http://stackoverflow.com/a/24336429 */
switch(err) {{
/* run-time and JIT compiler errors */
case 0: return "CL_SUCCESS";
case -1: return "CL_DEVICE_NOT_FOUND";
case -2: return "CL_DEVICE_NOT_AVAILABLE";
case -3: return "CL_COMPILER_NOT_AVAILABLE";
case -4: return "CL_MEM_OBJECT_ALLOCATION_FAILURE";
case -5: return "CL_OUT_OF_RESOURCES";
case -6: return "CL_OUT_OF_HOST_MEMORY";
case -7: return "CL_PROFILING_INFO_NOT_AVAILABLE";
case -8: return "CL_MEM_COPY_OVERLAP";
case -9: return "CL_IMAGE_FORMAT_MISMATCH";
case -10: return "CL_IMAGE_FORMAT_NOT_SUPPORTED";
case -11: return "CL_BUILD_PROGRAM_FAILURE";
case -12: return "CL_MAP_FAILURE";
case -13: return "CL_MISALIGNED_SUB_BUFFER_OFFSET";
case -14: return "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST";
case -15: return "CL_COMPILE_PROGRAM_FAILURE";
case -16: return "CL_LINKER_NOT_AVAILABLE";
case -17: return "CL_LINK_PROGRAM_FAILURE";
case -18: return "CL_DEVICE_PARTITION_FAILED";
case -19: return "CL_KERNEL_ARG_INFO_NOT_AVAILABLE";
/* compile-time errors */
case -30: return "CL_INVALID_VALUE";
case -31: return "CL_INVALID_DEVICE_TYPE";
case -32: return "CL_INVALID_PLATFORM";
case -33: return "CL_INVALID_DEVICE";
case -34: return "CL_INVALID_CONTEXT";
case -35: return "CL_INVALID_QUEUE_PROPERTIES";
case -36: return "CL_INVALID_COMMAND_QUEUE";
case -37: return "CL_INVALID_HOST_PTR";
case -38: return "CL_INVALID_MEM_OBJECT";
case -39: return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR";
case -40: return "CL_INVALID_IMAGE_SIZE";
case -41: return "CL_INVALID_SAMPLER";
case -42: return "CL_INVALID_BINARY";
case -43: return "CL_INVALID_BUILD_OPTIONS";
case -44: return "CL_INVALID_PROGRAM";
case -45: return "CL_INVALID_PROGRAM_EXECUTABLE";
case -46: return "CL_INVALID_KERNEL_NAME";
case -47: return "CL_INVALID_KERNEL_DEFINITION";
case -48: return "CL_INVALID_KERNEL";
case -49: return "CL_INVALID_ARG_INDEX";
case -50: return "CL_INVALID_ARG_VALUE";
case -51: return "CL_INVALID_ARG_SIZE";
case -52: return "CL_INVALID_KERNEL_ARGS";
case -53: return "CL_INVALID_WORK_DIMENSION";
case -54: return "CL_INVALID_WORK_GROUP_SIZE";
case -55: return "CL_INVALID_WORK_ITEM_SIZE";
case -56: return "CL_INVALID_GLOBAL_OFFSET";
case -57: return "CL_INVALID_EVENT_WAIT_LIST";
case -58: return "CL_INVALID_EVENT";
case -59: return "CL_INVALID_OPERATION";
case -60: return "CL_INVALID_GL_OBJECT";
case -61: return "CL_INVALID_BUFFER_SIZE";
case -62: return "CL_INVALID_MIP_LEVEL";
case -63: return "CL_INVALID_GLOBAL_WORK_SIZE";
case -64: return "CL_INVALID_PROPERTY";
case -65: return "CL_INVALID_IMAGE_DESCRIPTOR";
case -66: return "CL_INVALID_COMPILER_OPTIONS";
case -67: return "CL_INVALID_LINKER_OPTIONS";
case -68: return "CL_INVALID_DEVICE_PARTITION_COUNT";
/* extension errors */
case -1000: return "CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR";
case -1001: return "CL_PLATFORM_NOT_FOUND_KHR";
case -1002: return "CL_INVALID_D3D10_DEVICE_KHR";
case -1003: return "CL_INVALID_D3D10_RESOURCE_KHR";
case -1004: return "CL_D3D10_RESOURCE_ALREADY_ACQUIRED_KHR";
case -1005: return "CL_D3D10_RESOURCE_NOT_ACQUIRED_KHR";
default: return "Unknown OpenCL error";
}}
}}
void check_error(const char* api_call, cl_int err) {{
if(err != CL_SUCCESS) {{
fprintf(stderr, "%s %s\\n", api_call, clerror_string(err));
exit(1);
}}
}}
int help(char **argv) {{
printf("Usage: %s [-p <platform-id>] [-d <device-id>]\\n", argv[0]);
return 2;
}}
int main(int argc, char** argv) {{
int err;
int platform_id = PLATFORM_ID;
int device_id = DEVICE_ID;
const char *filename = NULL;
for (int i = 1; i < argc; i++) {{
if (!strcmp(argv[i], "-h") || !strcmp(argv[i], "--help"))
return help(argv);
else if (!strcmp(argv[i], "-f"))
filename = argv[++i];
else if (!strcmp(argv[i], "-p"))
platform_id = atoi(argv[++i]);
else if (!strcmp(argv[i], "-d"))
device_id = atoi(argv[++i]);
else
fprintf(stderr, "warning: unrecognized argument '%s'\\n", argv[i]);
}}
/* Optionally read kernel from file */
if (filename) {{
FILE *infile = fopen(filename, "rb");
if (infile == NULL) {{
fprintf(stderr, "fatal: Could not open '%s'\\n", filename);
return 3;
}}
fseek(infile, 0, SEEK_END);
long fsize = ftell(infile);
fseek(infile, 0, SEEK_SET);
char *buf = malloc(fsize + 1);
fread(buf, fsize, 1, infile);
fclose(infile);
buf[fsize] = 0;
fprintf(stderr, "read kernel from '%s'\\n", filename);
kernel_src = buf;
}}
cl_uint num_platforms;
cl_platform_id *platform_ids = (cl_platform_id*)malloc(sizeof(cl_platform_id) * (platform_id + 1));
err = clGetPlatformIDs(platform_id + 1, platform_ids, &num_platforms);
check_error("clGetPlatformIDs", err);
if (num_platforms <= platform_id) {{
fprintf(stderr, "Platform ID %d not found\\n", platform_id);
return 1;
}}
cl_platform_id cl_platform_id = platform_ids[platform_id];
char strbuf[256];
err = clGetPlatformInfo(cl_platform_id, CL_PLATFORM_NAME, sizeof(strbuf), strbuf, NULL);
check_error("clGetPlatformInfo", err);
fprintf(stderr, "[cldrive] Platform: %s\\n", strbuf);
cl_uint num_devices;
cl_device_id *device_ids = (cl_device_id*)malloc(sizeof(cl_device_id) * (device_id + 1));
err = clGetDeviceIDs(cl_platform_id, CL_DEVICE_TYPE_ALL, device_id + 1, device_ids, &num_devices);
check_error("clGetDeviceIDs", err);
if (num_devices <= device_id) {{
fprintf(stderr, "Device ID %d not found\\n", device_id);
return 1;
}}
cl_device_id cl_device_id = device_ids[device_id];
err = clGetDeviceInfo(cl_device_id, CL_DEVICE_NAME, sizeof(strbuf), strbuf, NULL);
check_error("clGetDeviceInfo", err);
fprintf(stderr, "[cldrive] Device: %s\\n", strbuf);
cl_context ctx = clCreateContext(NULL, 1, &cl_device_id, NULL, NULL, &err);
check_error("clCreateContext", err);
cl_command_queue queue = clCreateCommandQueue(ctx, cl_device_id, 0, &err);
check_error("clCreateCommandQueue", err);
fprintf(stderr, "[cldrive] OpenCL optimizations: {optimizations_on_off}\\n");
cl_program program = clCreateProgramWithSource(ctx, 1, (const char **) &kernel_src, NULL, &err);
check_error("clCreateProgramWithSource", err);
int build_err = clBuildProgram(program, 0, NULL, {clBuildProgram_opts}, NULL, NULL);
size_t log_size;
err = clGetProgramBuildInfo(program, cl_device_id, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
check_error("clGetProgramBuildInfo", err);
if (log_size > 2) {{
char* log = (char*)malloc(sizeof(char) * (log_size + 1));
err = clGetProgramBuildInfo(program, cl_device_id, CL_PROGRAM_BUILD_LOG, log_size, log, NULL);
check_error("clGetProgramBuildInfo", err);
fprintf(stderr, "%s", log);
}}
check_error("clBuildProgram", build_err);
"""
if not compile_only or (compile_only and create_kernel):
kernel_name_ = _args.GetKernelName(src)
c += f"""
cl_kernel kernels[128];
cl_uint num_kernels;
err = clCreateKernelsInProgram(program, 128, kernels, &num_kernels);
check_error("clCreateKernelsInProgram", err);
if (num_kernels != 1) {{
fprintf(stderr, "fatal: require 1 kernel, got %u\\n", num_kernels);
return 3;
}}
cl_kernel kernel = kernels[0];
char kernel_name[128];
err = clGetKernelInfo(kernel, CL_KERNEL_FUNCTION_NAME, 128, kernel_name, NULL);
check_error("clGetKernelInfo", err);
if (!filename && strcmp(kernel_name, "{kernel_name_}"))
fprintf(stderr, "fatal: expected kernel name \\\"{kernel_name_}\\\", got \\\"%s\\\"\\n", kernel_name);
fprintf(stderr, "[cldrive] Kernel: \\\"%s\\\"\\n", kernel_name);
"""
if not compile_only:
args = _args.GetKernelArguments(src)
setup_block, teardown_block, print_block = gen_data_blocks(args, inputs)
c += f"""
{setup_block}
const size_t lsize[3] = {{ {lsize.x}, {lsize.y}, {lsize.z} }};
const size_t gsize[3] = {{ {gsize.x}, {gsize.y}, {gsize.z} }};
err = clEnqueueNDRangeKernel(queue, kernel, 3, NULL, gsize, lsize, 0, NULL, NULL);
check_error("clEnqueueNDRangeKernel", err);
{teardown_block}
err = clFinish(queue);
check_error("clFinish", err);
{print_block}
clReleaseProgram(program);
clReleaseKernel(kernel);
clReleaseCommandQueue(queue);
clReleaseContext(ctx);
"""
# close out main():
c += f"""
fprintf(stderr, "done.\\n");
return 0;
}}
"""
return c
def HasVectorInputs(testcase: deepsmith_pb2.Testcase) -> bool:
"""Return whether any of the kernel arguments are vector types."""
for arg in args.GetKernelArguments(testcase.inputs['src']):
if arg.is_vector:
return True
return False
def MakeData(src: str, size: int, data_generator: Generator,
scalar_val: float = None) -> np.array:
"""Generate data for OpenCL kernels.
Creates a numpy array for each OpenCL argument, except arguments with the
'local' qualifier, since those are instantiated.
Returns:
The generated data as an np.array.
Raises:
TypeError: If any of the input arguments are of incorrect type.
ValueError: If any of the arguments cannot be interpreted.
Examples:
>>> MakeData("kernel void A(global int* a, const int b) {}", 3, Generator.ZEROS)
array([array([0, 0, 0], dtype=int32), array([3], dtype=int32)],
dtype=object)
>>> MakeData("kernel void A(global int* a, const int b) {}", 3, Generator.ONES)
array([array([1, 1, 1], dtype=int32), array([3], dtype=int32)],
dtype=object)
>>> MakeData("kernel void A(global int* a, const int b) {}", 3, Generator.ARANGE)
array([array([0, 1, 2], dtype=int32), array([3], dtype=int32)],
dtype=object)
Use `scalar_val` parameter to fix the value of scalar arguments:
>>> MakeData("kernel void A(global int* a, const int b) {}", 3, Generator.ARANGE, scalar_val=100)
array([array([0, 1, 2], dtype=int32), array([100], dtype=int32)],
dtype=object)
"""
# check the input types
err.assert_or_raise(isinstance(src, str), TypeError)
err.assert_or_raise(isinstance(data_generator, Generator), TypeError,
"invalid argument type for enum data_generator")
if scalar_val is None:
scalar_val = size
data = []
for arg in _args.GetKernelArguments(src):
if arg.address_space == "global" or arg.address_space == "constant":
argdata = data_generator(arg.numpy_type, size * arg.vector_width)
elif arg.address_space == "local":
# we don't need to generate data for local memory
continue
elif not arg.is_pointer:
# scalar values are still arrays, so e.g. 'float4' is an array of
# 4 floats. Each component of a scalar value is the flattened
# global size, e.g. with gsize (32,2,1), scalar arugments have the
# value 32 * 2 * 1 = 64.
argdata = np.array([scalar_val] * arg.vector_width).astype(arg.numpy_type)
else:
# argument is neither global or local, but is a pointer?
raise ValueError(f"unknown argument type '{arg}'")
data.append(argdata)
return np.array(data)
def main(argv):
"""Main entry point."""
if len(argv) > 1:
logging.warning("Unknown arguments: '{}'.".format(' '.join(argv[1:])))
if FLAGS.ls_env:
env.PrintOpenClEnvironments()
# Read kernel source.
src = sys.stdin.read()
# Parse inputs from strings.
gsize = driver.NDRange.FromString(FLAGS.gsize)
lsize = driver.NDRange.FromString(FLAGS.lsize)
data_generator = data.Generator.FromString(FLAGS.generator)
env_ = env.make_env(devtype=FLAGS.devtype,
platform=FLAGS.platform,
device=FLAGS.device)
if FLAGS.compile_only:
inputs = []
else:
inputs = data.MakeData(src=src,
size=FLAGS.size,
data_generator=data_generator,
scalar_val=FLAGS.scalar_val)
drive_args = {
"src": src,
"inputs": inputs,
"gsize": gsize,
"lsize": lsize,
"optimizations": not FLAGS.cl_opt,
"profiling": FLAGS.profiling,
"debug": FLAGS.debug,
"timeout": FLAGS.timeout,
}
if FLAGS.emit_c:
emit_c_args = {
"compile_only": FLAGS.compile_only,
"create_kernel": FLAGS.with_kernel,
}
print(cgen.emit_c(**drive_args, **emit_c_args))
else:
outputs = driver.DriveKernel(**drive_args, env=env_)
# Print result.
if FLAGS.binary:
d = pickle.dumps(outputs)
sys.stdout = io.TextIOWrapper(sys.stdout.detach(),
encoding='latin-1')
print(d.decode('latin-1'), end='', flush=True)
else:
np.set_printoptions(threshold=np.nan)
args_ = [
arg for arg in args.GetKernelArguments(src)
if not arg.address_space == 'local'
]
assert (len(args_) == len(outputs))
for arr, arg in zip(outputs, args_):
print(f"{arg.name}: {arr}")
def DriveKernel(env: _env.OpenCLEnvironment,
src: str,
inputs: np.array,
gsize: typing.Union[typing.Tuple[int, int, int], NDRange],
lsize: typing.Union[typing.Tuple[int, int, int], NDRange],
timeout: int = -1,
optimizations: bool = True,
profiling: bool = False,
debug: bool = False) -> np.array:
"""Drive an OpenCL kernel.
Executes an OpenCL kernel on the given environment, over the given inputs.
Execution is performed in a subprocess.
Args:
env: The OpenCL environment to run the kernel in.
src: The OpenCL kernel source.
inputs: The input data to the kernel.
optimizations: Whether to enable or disbale OpenCL compiler optimizations.
profiling: If true, print OpenCLevent times for data transfers and kernel
executions to stderr.
timeout: Cancel execution if it has not completed after this many seconds.
A value <= 0 means never time out.
debug: If true, silence the OpenCL compiler.
Returns:
A numpy array of the same shape as the inputs, with the values after
running the OpenCL kernel.
Raises:
ValueError: If input types are incorrect.
TypeError: If an input is of an incorrect type.
LogicError: If the input types do not match OpenCL kernel types.
PorcelainError: If the OpenCL subprocess exits with non-zero return code.
RuntimeError: If OpenCL program fails to build or run.
Examples:
A simple kernel which doubles its inputs:
>>> src = "kernel void A(global int* a) { a[get_global_id(0)] *= 2; }"
>>> inputs = [[1, 2, 3, 4, 5]]
>>> DriveKernel(env, src, inputs, gsize=(5,1,1), lsize=(1,1,1)) # doctest: +SKIP
array([[ 2, 4, 6, 8, 10]], dtype=int32)
"""
def Log(*args, **kwargs):
"""Log a message to stderr."""
if debug:
print("[cldrive] ", end="", file=sys.stderr)
print(*args, **kwargs, file=sys.stderr)
# Assert input types.
err.assert_or_raise(isinstance(env, _env.OpenCLEnvironment), ValueError,
"env argument is of incorrect type")
err.assert_or_raise(isinstance(src, str), ValueError,
"source is not a string")
# Validate global and local sizes.
err.assert_or_raise(len(gsize) == 3, TypeError)
err.assert_or_raise(len(lsize) == 3, TypeError)
gsize, lsize = NDRange(*gsize), NDRange(*lsize)
err.assert_or_raise(gsize.product >= 1, ValueError,
f"Scalar global size {gsize.product} must be >= 1")
err.assert_or_raise(lsize.product >= 1, ValueError,
f"Scalar local size {lsize.product} must be >= 1")
err.assert_or_raise(
gsize >= lsize, ValueError,
f"Global size {gsize} must be larger than local size {lsize}")
# Parse args in this process since we want to preserve the sueful exception
# type.
args = _args.GetKernelArguments(src)
# Check that the number of inputs is correct.
args_with_inputs = [
i for i, arg in enumerate(args) if not arg.address_space == 'local'
]
err.assert_or_raise(
len(args_with_inputs) == len(inputs), ValueError,
"Kernel expects {} inputs, but {} were provided".format(
len(args_with_inputs), len(inputs)))
# All inputs must have some length.
for i, x in enumerate(inputs):
err.assert_or_raise(len(x), ValueError, f"Input {i} has size zero")
# Copy inputs into the expected data types.
data = np.array(
[np.array(d).astype(a.numpy_type) for d, a in zip(inputs, args)])
job = {
"env": env,
"src": src,
"args": args,
"data": data,
"gsize": gsize,
"lsize": lsize,
"optimizations": optimizations,
"profiling": profiling
}
with NamedTemporaryFile('rb+', prefix='cldrive-',
suffix='.job') as tmp_file:
porcelain_job_file = tmp_file.name
# Write job file.
pickle.dump(job, tmp_file)
tmp_file.flush()
# Enforce timeout using sigkill.
if timeout > 0:
cli = ["timeout", "--signal=9", str(int(timeout))]
else:
cli = []
cli += [sys.executable, __file__, porcelain_job_file]
cli_str = " ".join(cli)
Log("Porcelain invocation:", cli_str)
# Fork and run.
process = Popen(cli, stdout=PIPE, stderr=PIPE)
stdout, stderr = process.communicate()
status = process.returncode
if debug:
print(stdout.decode('utf-8').strip(), file=sys.stderr)
print(stderr.decode('utf-8').strip(), file=sys.stderr)
elif profiling:
# Print profiling output when not in debug mode.
for line in stderr.decode('utf-8').split('\n'):
if re.match(r'\[cldrive\] .+ time: [0-9]+\.[0-9]+ ms', line):
print(line, file=sys.stderr)
Log(f"Porcelain return code: {status}")
# Test for non-zero exit codes. The porcelain subprocess catches exceptions
# and completes gracefully, so a non-zero return code is indicative of a
# more serious problem.
#
# FIXME: I'm seeing a number of SIGABRT return codes which I can't explain.
# However, ignoring them seems to not cause a problem ...
if status != 0 and status != -Signals['SIGABRT'].value:
# A negative return code means a signal. Try and convert the value into a
# signal name.
with suppress(ValueError):
status = Signals(-status).name
if status == "SIGKILL":
raise TimeoutError(timeout)
else:
raise PorcelainError(status)
# Read result.
tmp_file.seek(0)
return_value = pickle.load(tmp_file)
outputs = return_value["outputs"]
error = return_value["err"]
if error: # Porcelain raised an exception, re-raise it.
raise error
else:
return outputs
def test_GetKernelArguments_no_args():
"""Test that no arguments returned for kernel with no args."""
assert len(args.GetKernelArguments("kernel void A() {}")) == 0
def test_GetKernelArguments_no_qualifiers():
"""Test that error is raised if argument has no address space qualifier."""
with pytest.raises(args.OpenCLValueError) as e_ctx:
args.GetKernelArguments("kernel void A(float* a) {}")
assert "Pointer argument 'float *a' has no address space qualifier" == str(
e_ctx.value)
def test_GetKernelArguments_local_global_qualified():
"""Test that error is raised if address space is invalid."""
with pytest.raises(args.OpenCLValueError) as e_ctx:
args.GetKernelArguments("kernel void A(global local int* a) {}")
assert ("Pointer argument 'global local int *a' has multiple "
'address space qualifiers') == str(e_ctx.value)
def test_GetKernelArguments_struct_not_supported():
"""Test that error is raised if type is not supported."""
with pytest.raises(ValueError) as e_ctx:
args.GetKernelArguments("struct C; kernel void A(struct C a) {}")
assert "Unsupported data type for argument: 'a'" == str(e_ctx.value)
def test_GetKernelArguments_declared_but_not_defined():
"""Test that error is raised if kernel declared but not defined."""
with pytest.raises(args.NoKernelError):
args.GetKernelArguments("kernel void A();")
def test_GetKernelArguments_no_definition():
"""Test that error is raised if no kernel defined."""
with pytest.raises(args.NoKernelError):
args.GetKernelArguments("")
def test_GetKernelArguments_hello_world():
"""Simple hello world argument type test."""
args_ = args.GetKernelArguments("kernel void a(global float* a) {}")
assert args_[0].typename == 'float'
