def test_GetKernelArguments_multiple_kernels():
"""Test that error is raised if no kernel defined."""
with test.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 GetKernelArguments(kernel: str):
try:
# Extract everything up to the function body, and use an empty function
# body for parsing. This means that errors that are in the function body
# will not cause this to fail. E.g. given kernel:
#
# kernel void A(const int a, global float* b) {
# b[0] += a;
# }
#
# This will parse:
#
# kernel void A(const int a, global float* b) {}
kernel_declaration = kernel[:kernel.index("{")] + "{}"
return args.GetKernelArguments(kernel_declaration)
except ValueError as e:
app.Error("Failure processing kernel: '%s'", kernel)
raise e
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 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 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 test.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 test.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 test.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 test.Raises(args.NoKernelError):
args.GetKernelArguments("kernel void A();")
def test_GetKernelArguments_no_definition():
"""Test that error is raised if no kernel defined."""
with test.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"
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 = (char *)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 main(argv):
"""Main entry point."""
if len(argv) > 1:
app.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}")