Port build_module.py to C++ #667

alex-weaver · 2017-11-23T11:33:10Z

I've ported the functions from build_module.py over to C++ to allow compiling from C++; this required also implementing equivalents of Target and BuildConfig in C++. I've also added EXPORT directives to a minimal set of functions in schedule.h to allow implementing the Getting Started example. I would imagine in future most of schedule.h will need EXPORT directives to allow specifying more elaborate schedules from C++.

This implementation can be verified with the following port of the Getting Started example:

#include <tvm\operation.h>
#include <tvm\compilation.h>


using namespace tvm;
using namespace tvm::compilation;

int main()
{
	auto n = tvm::Variable::make(tvm::Int(32), "n");
	tvm::Array<tvm::Expr> shape;
	shape.push_back(n);

	auto A = tvm::placeholder(shape, tvm::Float(32), "A");
	auto B = tvm::placeholder(shape, tvm::Float(32), "B");

	auto C = tvm::compute(A->shape, [&A, &B](tvm::Expr i) {
		return A[i] + B[i];
	}, "C");

	auto s = tvm::create_schedule({ C->op });
	

	auto cAxis = GetAttrHandle<Array<IterVar>>(C->op, "axis");

	tvm::IterVar bx, tx;
	s[C].split(cAxis[0], 64, &bx, &tx);

	s[C].bind(bx, tvm::thread_axis(tvm::Range(), "blockIdx.x"));
	s[C].bind(tx, tvm::thread_axis(tvm::Range(), "threadIdx.x"));
	

	auto args = Array<Tensor>({ A, B, C });
	std::unordered_map<Tensor, Buffer> binds;

	BuildConfig config;
	auto target = target_cuda();
	auto targetHost = default_target_host(target);

	auto lowered = Lower(s, args, "func", binds, config);
	auto module = BuildModule(Array<LoweredFunc>({ lowered }), target, targetHost, config);

	auto dev_module = module->imports()[0];
	std::cout << dev_module->GetSource("");

    return 0;
}

tqchen · 2017-11-23T19:22:02Z

Thanks for starting this. Since the C++ api headers is quite serious change, there are a few things that we need to fix in this PR. Here are some guidelines

Code style
- Use stl_case for variables and structure members
- Use CamelCase for internal functions
- Mark inline for functions in header files
  - Aoid doing functions in the header files unless they are strictly short.
User facing API
- Use stl_case for the user facing functions(put them under tvm namespace)
- This include lower and build
- Make the function signature as consistent as possible with python counterparts, as eventually we can expose these functions

tqchen · 2017-11-23T19:22:26Z

include/tvm/compilation.h

+* \file compilation.h
+* \brief Functions for compiling ops.
+*/
+#ifndef TVM_COMPILATION_H_


use the same name tvm/build_module.h

tqchen · 2017-11-23T19:22:43Z

include/tvm/compilation.h

+
+namespace tvm {
+
+namespace compilation {


leave user facing API and data structure under root namespace.

tqchen · 2017-11-23T19:23:06Z

include/tvm/compilation.h

+    std::unordered_set<std::string> options;
+
+
+    Target(std::string targetName, DLDeviceType deviceType, int max_num_threads,


add a static function to construct from target string

tqchen · 2017-11-23T19:23:16Z

include/tvm/compilation.h

+*/
+struct Target {
+    /*! \brief The name of the target device */
+    std::string targetName;


use stl_case for variable name

tqchen · 2017-11-23T19:23:44Z

include/tvm/compilation.h

+    /*! \brief The type of the target device */
+    DLDeviceType deviceType;
+    /*! \brief The maximum threads that a schedule should use for this device */
+    int max_num_threads;


set default value of the optional argument here

tqchen · 2017-11-23T19:28:46Z

include/tvm/compilation.h

+
+
+/*! \brief Convenience function for getting attributes */
+TVMValue GetAttr(NodeRef node, std::string attrName) {


This function do not belong to here. Instead of using GetAttr, use

node.as<Expr>->type() for the type

tqchen · 2017-11-23T19:29:15Z

include/tvm/compilation.h

+
+/*! \brief Convenience function for getting handle attributes */
+template<typename T>
+T GetAttrHandle(NodeRef node, std::string attrName) {


use the node.as<> instead of C API mechanism to get the children

tqchen · 2017-11-23T19:29:25Z

include/tvm/compilation.h

+* \param config The build configuration.
+* \return The built Stmt.
+*/
+EXPORT Stmt BuildStmt(Schedule sch, Array<Tensor> args, std::unordered_map<Tensor, Buffer> binds,


This is a private function that should belong to cc file

tqchen · 2017-11-23T19:30:03Z

include/tvm/compilation.h

+* \param config The build configuration.
+* \return The lowered function.
+*/
+EXPORT LoweredFunc Lower(Schedule sch, Array<Tensor> args, std::string name,


Lower->lower as it is user facing, make API function consistent with python API. Pass by const reference when possible if the function do not move the content

returns Array

tqchen · 2017-11-23T19:30:39Z

include/tvm/compilation.h

+* \param config The build configuration.
+* \return The built module.
+*/
+EXPORT runtime::Module BuildModule(Array<LoweredFunc> funcs, const Target& target,


BuildModule->build

alex-weaver · 2017-11-24T13:11:22Z

Ok hopefully that commit resolves the style issues - naming conventions should now be observed and the API surface in the header should be significantly reduced. There's a couple of things I'm not sure about though:

You mentioned lower() should return array - the python version of lower() does not return array - is this correct?
You requested a static function to construct a Target from a string - is this intended to just switch on the name and call the appropriate function in target:: or should it parse out options too?

tqchen · 2017-11-24T18:23:30Z

The main reason is that the python version of build is be able to accept both schedule/ array of LoweredFunc and LoweredFunc.
- We want result of lower to be directly feedable to the build function, so as a compromise here, we can have lower return a list of array for now.
It should be consistent with the Target parsing behavior on python side

lower now returns array

alex-weaver · 2017-11-24T19:37:24Z

Ok lower() now returns array, and there is a function Target::create which should be consistent with create() in target.py

tqchen

Thanks for being patient with this, here is another set of review comments

tqchen · 2017-11-27T19:30:23Z

include/tvm/build_module.h

+
+/*!
+* \brief Container for target device information.
+* Use target_llvm, target_cuda etc functions instead of constructing directly.


update comment to target::llvm

tqchen · 2017-11-27T19:30:30Z

include/tvm/build_module.h

+#define TVM_BUILD_MODULE_H_
+
+#include <string>
+#include "./tvm/c_dsl_api.h"


this is no longer needed

tqchen · 2017-11-27T19:30:41Z

include/tvm/build_module.h

+    EXPORT static Target create(const std::string& target_str);
+};
+
+namespace target {


document the namespace

tqchen · 2017-11-27T19:31:55Z

include/tvm/build_module.h

+
+    Target(const std::string& target_name, DLDeviceType device_type, int max_num_threads,
+        int thread_warp_size, const std::unordered_set<std::string>& keys,
+        const std::unordered_set<std::string>& options) {


use member initializers for class members.

tqchen · 2017-11-27T19:32:20Z

include/tvm/build_module.h

+
+namespace tvm {
+
+/*!


general style issue: Google C style requires 2-space indentation.

tqchen · 2017-11-27T19:37:54Z