inline kernel_call bluestein_pad_kernel(const cl::CommandQueue &queue, size_t n, size_t m, const cl::Buffer &in, const cl::Buffer &out) { std::ostringstream o; kernel_common<T>(o, qdev(queue)); o << "real2_t conj(real2_t v) {\n" << " return (real2_t)(v.x, -v.y);\n" << "}\n"; o << "__kernel void bluestein_pad_kernel(" << "__global real2_t *input, __global real2_t *output, uint n, uint m) {\n" << " const size_t x = get_global_id(0);\n" << " if(x < n || m - x < n)\n" << " output[x] = conj(input[min(x, m - x)]);\n" << " else\n" << " output[x] = (real2_t)(0,0);\n" << "}\n"; auto program = build_sources(qctx(queue), o.str()); cl::Kernel kernel(program, "bluestein_pad_kernel"); kernel.setArg(0, in); kernel.setArg(1, out); kernel.setArg(2, static_cast<cl_uint>(n)); kernel.setArg(3, static_cast<cl_uint>(m)); std::ostringstream desc; desc << "bluestein_pad_kernel{n=" << n << ", m=" << m << "}"; return kernel_call(true, desc.str(), program, kernel, cl::NDRange(m), cl::NullRange); }
inline kernel_call bluestein_mul(const cl::CommandQueue &queue, size_t n, size_t batch, const cl::Buffer &data, const cl::Buffer &exp, const cl::Buffer &out) { std::ostringstream o; kernel_common<T>(o, qdev(queue)); mul_code(o, false); o << "__kernel void bluestein_mul(" << "__global const real2_t *data, __global const real2_t *exp, __global real2_t *output, uint stride) {\n" << " const size_t x = get_global_id(0), y = get_global_id(1);\n" << " if(x < stride) {\n" << " const size_t off = x + stride * y;" << " output[off] = mul(data[off], exp[x]);\n" << " }\n" << "}\n"; auto program = build_sources(qctx(queue), o.str()); cl::Kernel kernel(program, "bluestein_mul"); kernel.setArg(0, data); kernel.setArg(1, exp); kernel.setArg(2, out); kernel.setArg(3, static_cast<cl_uint>(n)); const size_t wg = kernel.getWorkGroupInfo<CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE>(qdev(queue)); const size_t threads = alignup(n, wg); std::ostringstream desc; desc << "bluestein_mul{n=" << n << "(" << threads << "), wg=" << wg << ", batch=" << batch << "}"; return kernel_call(false, desc.str(), program, kernel, cl::NDRange(threads, batch), cl::NDRange(wg, 1)); }
inline kernel_call radix_kernel(bool once, const cl::CommandQueue &queue, size_t n, size_t batch, bool invert, pow radix, size_t p, const cl::Buffer &in, const cl::Buffer &out) { std::ostringstream o; o << std::setprecision(25); const auto device = qdev(queue); kernel_common<T>(o, device); mul_code(o, invert); twiddle_code<T>(o); const size_t m = n / radix.value; kernel_radix<T>(o, radix, invert); auto program = build_sources(qctx(queue), o.str(), "-cl-mad-enable -cl-fast-relaxed-math"); cl::Kernel kernel(program, "radix"); kernel.setArg(0, in); kernel.setArg(1, out); kernel.setArg(2, static_cast<cl_uint>(p)); kernel.setArg(3, static_cast<cl_uint>(m)); const size_t wg_mul = kernel.getWorkGroupInfo<CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE>(device); //const size_t max_cu = device.getInfo<CL_DEVICE_MAX_COMPUTE_UNITS>(); //const size_t max_wg = device.getInfo<CL_DEVICE_MAX_WORK_GROUP_SIZE>(); size_t wg = wg_mul; //while(wg * max_cu < max_wg) wg += wg_mul; //wg -= wg_mul; const size_t threads = alignup(m, wg); std::ostringstream desc; desc << "dft{r=" << radix << ", p=" << p << ", n=" << n << ", batch=" << batch << ", threads=" << m << "(" << threads << "), wg=" << wg << "}"; return kernel_call(once, desc.str(), program, kernel, cl::NDRange(threads, batch), cl::NDRange(wg, 1)); }
/// Constructor. Creates a cl::Kernel instance from source. kernel(const cl::CommandQueue &queue, const std::string &src, const std::string &name, std::function<size_t(size_t)> smem ) : argpos(0), K(build_sources(queue, src), name.c_str()) { config(queue, smem); }
inline kernel_call transpose_kernel(const cl::CommandQueue &queue, size_t width, size_t height, const cl::Buffer &in, const cl::Buffer &out) { std::ostringstream o; const auto dev = qdev(queue); kernel_common<T>(o, dev); // determine max block size to fit into local memory/workgroup size_t block_size = 128; { const auto local_size = dev.getInfo<CL_DEVICE_LOCAL_MEM_SIZE>(); const auto workgroup = dev.getInfo<CL_DEVICE_MAX_WORK_GROUP_SIZE>(); while(block_size * block_size * sizeof(T) * 2 > local_size) block_size /= 2; while(block_size * block_size > workgroup) block_size /= 2; } // from NVIDIA SDK. o << "__kernel void transpose(" << "__global const real2_t *input, __global real2_t *output, uint width, uint height) {\n" << " const size_t " << " global_x = get_global_id(0), global_y = get_global_id(1),\n" << " local_x = get_local_id(0), local_y = get_local_id(1),\n" << " group_x = get_group_id(0), group_y = get_group_id(1),\n" << " block_size = " << block_size << ",\n" << " target_x = local_y + group_y * block_size,\n" << " target_y = local_x + group_x * block_size;\n" << " const bool range = global_x < width && global_y < height;\n" // local memory << " __local real2_t block[" << (block_size * block_size) << "];\n" // copy from input to local memory << " if(range)\n" << " block[local_x + local_y * block_size] = input[global_x + global_y * width];\n" // wait until the whole block is filled << " barrier(CLK_LOCAL_MEM_FENCE);\n" // transpose local block to target << " if(range)\n" << " output[target_x + target_y * height] = block[local_x + local_y * block_size];\n" << "}\n"; auto program = build_sources(qctx(queue), o.str()); cl::Kernel kernel(program, "transpose"); kernel.setArg(0, in); kernel.setArg(1, out); kernel.setArg(2, static_cast<cl_uint>(width)); kernel.setArg(3, static_cast<cl_uint>(height)); // range multiple of wg size, last block maybe not completely filled. size_t r_w = alignup(width, block_size); size_t r_h = alignup(height, block_size); std::ostringstream desc; desc << "transpose{" << "w=" << width << "(" << r_w << "), " << "h=" << height << "(" << r_h << "), " << "bs=" << block_size << "}"; return kernel_call(false, desc.str(), program, kernel, cl::NDRange(r_w, r_h), cl::NDRange(block_size, block_size)); }
/// Constructor. Creates a backend::kernel instance from source. kernel(const boost::compute::command_queue &queue, const std::string &src, const std::string &name, std::function<size_t(size_t)> smem, const std::string &options = "" ) : argpos(0), K(build_sources(queue, src, options), name) { config(queue, smem); }
/// Constructor. Creates a backend::kernel instance from source. kernel(const command_queue &queue, const std::string &src, const std::string &name, std::function<size_t(size_t)> smem, const std::string &options = "" ) : ctx(queue.context()), P(build_sources(queue, src, options)), smem(0) { cuda_check( cuModuleGetFunction(&K, P.raw(), name.c_str()) ); config(queue, smem); }
/// Constructor. Creates a cl::Kernel instance from source. kernel(const cl::CommandQueue &queue, const std::string &src, const std::string &name, size_t smem_per_thread = 0 ) : argpos(0), K(build_sources(queue, src), name.c_str()) { config(queue, [smem_per_thread](size_t wgs){ return wgs * smem_per_thread; }); }
/// Constructor. Creates a backend::kernel instance from source. kernel(const boost::compute::command_queue &queue, const std::string &src, const std::string &name, size_t smem_per_thread = 0, const std::string &options = "" ) : argpos(0), K(build_sources(queue, src, options), name) { config(queue, [smem_per_thread](size_t wgs){ return wgs * smem_per_thread; }); }
/// Constructor. Creates a backend::kernel instance from source. kernel(const command_queue &queue, const std::string &src, const std::string &name, size_t smem_per_thread = 0, const std::string &options = "" ) : ctx(queue.context()), P(build_sources(queue, src, options)), smem(0) { cuda_check( cuModuleGetFunction(&K, P.raw(), name.c_str()) ); config(queue, [smem_per_thread](size_t wgs){ return wgs * smem_per_thread; }); }
inline kernel_call bluestein_mul_in(const cl::CommandQueue &queue, bool inverse, size_t batch, size_t radix, size_t p, size_t threads, size_t stride, const cl::Buffer &data, const cl::Buffer &exp, const cl::Buffer &out) { std::ostringstream o; kernel_common<T>(o, qdev(queue)); mul_code(o, false); twiddle_code<T>(o); o << "__kernel void bluestein_mul_in(" << "__global const real2_t *data, __global const real2_t *exp, __global real2_t *output, " << "uint radix, uint p, uint out_stride) {\n" << " const size_t\n" << " thread = get_global_id(0), threads = get_global_size(0),\n" << " batch = get_global_id(1),\n" << " element = get_global_id(2);\n" << " if(element < out_stride) {\n" << " const size_t\n" << " in_off = thread + batch * radix * threads + element * threads,\n" << " out_off = thread * out_stride + batch * out_stride * threads + element;\n" << " if(element < radix) {\n" << " real2_t w = exp[element];" << " if(p != 1) {\n" << " const int sign = " << (inverse ? "+1" : "-1") << ";\n" << " ulong a = (ulong)element * (thread % p);\n" << " ulong b = (ulong)radix * p;\n" << " real2_t t = twiddle(2 * sign * M_PI * (a % (2 * b)) / b);\n" << " w = mul(w, t);\n" << " }\n" << " output[out_off] = mul(data[in_off], w);\n" << " } else\n" << " output[out_off] = (real2_t)(0,0);" << " }\n" << "}\n"; auto program = build_sources(qctx(queue), o.str()); cl::Kernel kernel(program, "bluestein_mul_in"); kernel.setArg(0, data); kernel.setArg(1, exp); kernel.setArg(2, out); kernel.setArg(3, static_cast<cl_uint>(radix)); kernel.setArg(4, static_cast<cl_uint>(p)); kernel.setArg(5, static_cast<cl_uint>(stride)); const size_t wg = kernel.getWorkGroupInfo<CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE>(qdev(queue)); const size_t stride_pad = alignup(stride, wg); std::ostringstream desc; desc << "bluestein_mul_in{batch=" << batch << ", radix=" << radix << ", p=" << p << ", threads=" << threads << ", stride=" << stride << "(" << stride_pad << "), wg=" << wg << "}"; return kernel_call(false, desc.str(), program, kernel, cl::NDRange(threads, batch, stride_pad), cl::NDRange(1, 1, wg)); }
inline kernel_call bluestein_twiddle(const cl::CommandQueue &queue, size_t n, bool inverse, const cl::Buffer &out) { std::ostringstream o; kernel_common<T>(o, qdev(queue)); twiddle_code<T>(o); o << "__kernel void bluestein_twiddle(__global real2_t *output) {\n" << " const size_t x = get_global_id(0), n = get_global_size(0);\n" << " const int sign = " << (inverse ? "+1" : "-1") << ";\n" << " const size_t xx = ((ulong)x * x) % (2 * n);\n" << " output[x] = twiddle(sign * M_PI * xx / n);\n" << "}\n"; auto program = build_sources(qctx(queue), o.str()); cl::Kernel kernel(program, "bluestein_twiddle"); kernel.setArg(0, out); std::ostringstream desc; desc << "bluestein_twiddle{n=" << n << ", inverse=" << inverse << "}"; return kernel_call(true, desc.str(), program, kernel, cl::NDRange(n), cl::NullRange); }
inline kernel_call bluestein_mul_out(const cl::CommandQueue &queue, size_t batch, size_t p, size_t radix, size_t threads, size_t stride, const cl::Buffer &data, const cl::Buffer &exp, const cl::Buffer &out) { std::ostringstream o; kernel_common<T>(o, qdev(queue)); mul_code(o, false); o << "__kernel void bluestein_mul_out(" << "__global const real2_t *data, __global const real2_t *exp, __global real2_t *output, " << "real_t div, uint p, uint in_stride, uint radix) {\n" << " const size_t\n" << " i = get_global_id(0), threads = get_global_size(0),\n" << " b = get_global_id(1),\n" << " l = get_global_id(2);\n" << " if(l < radix) {\n" << " const size_t\n" << " k = i % p,\n" << " j = k + (i - k) * radix,\n" << " in_off = i * in_stride + b * in_stride * threads + l,\n" << " out_off = j + b * threads * radix + l * p;\n" << " output[out_off] = mul(data[in_off] * div, exp[l]);\n" << " }\n" << "}\n"; auto program = build_sources(qctx(queue), o.str()); cl::Kernel kernel(program, "bluestein_mul_out"); kernel.setArg(0, data); kernel.setArg(1, exp); kernel.setArg(2, out); kernel.setArg<T>(3, static_cast<T>(1) / stride); kernel.setArg(4, static_cast<cl_uint>(p)); kernel.setArg(5, static_cast<cl_uint>(stride)); kernel.setArg(6, static_cast<cl_uint>(radix)); const size_t wg = kernel.getWorkGroupInfo<CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE>(qdev(queue)); const size_t radix_pad = alignup(radix, wg); std::ostringstream desc; desc << "bluestein_mul_out{r=" << radix << "(" << radix_pad << "), wg=" << wg << ", batch=" << batch << ", p=" << p << ", thr=" << threads << ", stride=" << stride << "}"; return kernel_call(false, desc.str(), program, kernel, cl::NDRange(threads, batch, radix_pad), cl::NDRange(1, 1, wg)); }
typename std::enable_if< !boost::proto::matches< typename boost::proto::result_of::as_expr<ExprTuple>::type, multivector_full_grammar >::value #if !defined(_MSC_VER) || _MSC_VER >= 1700 && N == std::tuple_size<ExprTuple>::value #endif , const multivector& >::type operator=(const ExprTuple &expr) { #endif static kernel_cache cache; const std::vector<cl::CommandQueue> &queue = vec[0]->queue_list(); for(uint d = 0; d < queue.size(); d++) { cl::Context context = qctx(queue[d]); cl::Device device = qdev(queue[d]); auto kernel = cache.find( context() ); if (kernel == cache.end()) { std::ostringstream source; source << standard_kernel_header(device); { get_header f(source); for_each<0>(expr, f); } source << "kernel void multi_expr_tuple(\n" "\t" << type_name<size_t>() << " n"; for(uint i = 1; i <= N; i++) source << ",\n\tglobal " << type_name<T>() << " *res_" << i; { get_params f(source); for_each<0>(expr, f); } source << "\n)\n{\n"; if ( is_cpu(device) ) { source << "\tsize_t chunk_size = (n + get_global_size(0) - 1) / get_global_size(0);\n" "\tsize_t chunk_start = get_global_id(0) * chunk_size;\n" "\tsize_t chunk_end = min(n, chunk_start + chunk_size);\n" "\tfor(size_t idx = chunk_start; idx < chunk_end; ++idx) {\n"; } else { source << "\tfor(size_t idx = get_global_id(0); idx < n; idx += get_global_size(0)) {\n"; } { get_expressions f(source); for_each<0>(expr, f); } source << "\n"; for(uint i = 1; i <= N; i++) source << "\t\tres_" << i << "[idx] = buf_" << i << ";\n"; source << "\t}\n}\n"; auto program = build_sources(context, source.str()); cl::Kernel krn(program, "multi_expr_tuple"); size_t wgs = kernel_workgroup_size(krn, device); kernel = cache.insert(std::make_pair( context(), kernel_cache_entry(krn, wgs) )).first; } if (size_t psize = vec[0]->part_size(d)) { size_t w_size = kernel->second.wgsize; size_t g_size = num_workgroups(device) * w_size; uint pos = 0; kernel->second.kernel.setArg(pos++, psize); for(uint i = 0; i < N; i++) kernel->second.kernel.setArg(pos++, (*vec[i])(d)); { set_arguments f(kernel->second.kernel, d, pos, vec[0]->part_start(d)); for_each<0>(expr, f); } queue[d].enqueueNDRangeKernel( kernel->second.kernel, cl::NullRange, g_size, w_size ); } } return *this; }
typename std::enable_if< boost::proto::matches< typename boost::proto::result_of::as_expr<Expr>::type, multivector_expr_grammar >::value, const multivector& >::type operator=(const Expr& expr) { static kernel_cache cache; const std::vector<cl::CommandQueue> &queue = vec[0]->queue_list(); // If any device in context is CPU, then do not fuse the kernel, // but assign components individually. if (std::any_of(queue.begin(), queue.end(), [](const cl::CommandQueue &q) { return is_cpu(qdev(q)); })) { assign_subexpressions<0, N>(boost::proto::as_child(expr)); return *this; } for(uint d = 0; d < queue.size(); d++) { cl::Context context = qctx(queue[d]); cl::Device device = qdev(queue[d]); auto kernel = cache.find( context() ); if (kernel == cache.end()) { std::ostringstream kernel_name; kernel_name << "multi_"; vector_name_context name_ctx(kernel_name); boost::proto::eval(boost::proto::as_child(expr), name_ctx); std::ostringstream source; source << standard_kernel_header(device); extract_user_functions()( boost::proto::as_child(expr), declare_user_function(source) ); source << "kernel void " << kernel_name.str() << "(\n\t" << type_name<size_t>() << " n"; for(size_t i = 0; i < N; ) source << ",\n\tglobal " << type_name<T>() << " *res_" << ++i; build_param_list<N>(boost::proto::as_child(expr), source); source << "\n)\n{\n" "\tfor(size_t idx = get_global_id(0); idx < n; idx += get_global_size(0)) {\n"; build_expr_list(boost::proto::as_child(expr), source); source << "\t}\n}\n"; auto program = build_sources(context, source.str()); cl::Kernel krn(program, kernel_name.str().c_str()); size_t wgs = kernel_workgroup_size(krn, device); kernel = cache.insert(std::make_pair( context(), kernel_cache_entry(krn, wgs) )).first; } if (size_t psize = vec[0]->part_size(d)) { size_t w_size = kernel->second.wgsize; size_t g_size = num_workgroups(device) * w_size; uint pos = 0; kernel->second.kernel.setArg(pos++, psize); for(uint i = 0; i < N; i++) kernel->second.kernel.setArg(pos++, vec[i]->operator()(d)); set_kernel_args<N>( boost::proto::as_child(expr), kernel->second.kernel, d, pos, vec[0]->part_start(d) ); queue[d].enqueueNDRangeKernel( kernel->second.kernel, cl::NullRange, g_size, w_size ); } } return *this; }