114    auto &kern = obj(d_kern);  in clGetKernelInfo()  local
118       buf.as_string() = kern.name();  in clGetKernelInfo()
122       buf.as_scalar<cl_uint>() = kern.args().size();  in clGetKernelInfo()
126       buf.as_scalar<cl_uint>() = kern.ref_count();  in clGetKernelInfo()
130       buf.as_scalar<cl_context>() = desc(kern.program().context());  in clGetKernelInfo()
134       buf.as_scalar<cl_program>() = desc(kern.program());  in clGetKernelInfo()
138       buf.as_string() = find(name_equals(kern.name()), kern.program().symbols()).attributes;  in clGetKernelInfo()
156    auto &kern = obj(d_kern);  in clGetKernelWorkGroupInfo()  local
157    auto &dev = (d_dev ? *pobj(d_dev) : unique(kern.program().devices()));  in clGetKernelWorkGroupInfo()
159    if (!count(dev, kern.program().devices()))  in clGetKernelWorkGroupInfo()
168       buf.as_vector<size_t>() = kern.required_block_size();  in clGetKernelWorkGroupInfo()
172       buf.as_scalar<cl_ulong>() = kern.mem_local();  in clGetKernelWorkGroupInfo()
180       buf.as_scalar<cl_ulong>() = kern.mem_private();  in clGetKernelWorkGroupInfo()
246    validate_common(const command_queue &q, kernel &kern,  in validate_common()  argument
248       if (kern.program().context() != q.context() ||  in validate_common()
256             }, kern.args()))  in validate_common()
261       auto &b = kern.program().build(q.device()).bin;  in validate_common()
287    validate_block_size(const command_queue &q, const kernel &kern,  in validate_block_size()  argument
309          return kern.optimal_block_size(q, grid_size);  in validate_block_size()
321    auto &kern = obj(d_kern);  in clEnqueueNDRangeKernel()  local
325    auto block_size = validate_block_size(q, kern, dims,  in clEnqueueNDRangeKernel()
328    validate_common(q, kern, deps);  in clEnqueueNDRangeKernel()
332       [=, &kern, &q](event &) {  in clEnqueueNDRangeKernel()
333          kern.launch(q, grid_offset, grid_size, block_size);  in clEnqueueNDRangeKernel()
348    auto &kern = obj(d_kern);  in clEnqueueTask()  local
351    validate_common(q, kern, deps);  in clEnqueueTask()
355       [=, &kern, &q](event &) {  in clEnqueueTask()
356          kern.launch(q, { 0 }, { 1 }, { 1 });  in clEnqueueTask()
400    auto &kern = obj(d_kern);  in clSetKernelExecInfo()  local
403                                       kern.program().context().devices());  in clSetKernelExecInfo()