Skip to content
This repository has been archived by the owner on Jan 24, 2024. It is now read-only.

Commit

Permalink
Merge branch 'workspace' into add_op_test
Browse files Browse the repository at this point in the history
  • Loading branch information
@Xreki
Xreki committed Nov 17, 2021
2 parents 6baf60b + 8a77d5e commit 3221e6a
Show file tree
Hide file tree
Showing 43 changed files with 1,270 additions and 154 deletions.
4 changes: 2 additions & 2 deletions cinn/backends/codegen_cuda_dev_test.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -468,7 +468,7 @@ void __launch_bounds__(224) schedule_conv2d_0(const float* __restrict__ X, const
std::string trimed_source_target = utils::Trim(source_target);
int start_target = trimed_source_target.find("blockIdx");
int start_source = source_code.find("blockIdx");
ASSERT_EQ(trimed_source_target.substr(start_target), source_code.substr(start_source));
// ASSERT_EQ(trimed_source_target.substr(start_target), source_code.substr(start_source));
using runtime::cuda::CUDAModule;

backends::NVRTC_Compiler compiler;
Expand Down Expand Up @@ -615,7 +615,7 @@ void __launch_bounds__(128) schedule_conv2d_1(const float* __restrict__ X, const
std::string trimed_source_target = utils::Trim(source_target);
int start_target = trimed_source_target.find("blockIdx");
int start_source = source_code.find("blockIdx");
ASSERT_EQ(trimed_source_target.substr(start_target), source_code.substr(start_source));
// ASSERT_EQ(trimed_source_target.substr(start_target), source_code.substr(start_source));
using runtime::cuda::CUDAModule;

backends::NVRTC_Compiler compiler;
Expand Down
3 changes: 3 additions & 0 deletions cinn/backends/llvm/codegen_llvm.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -635,11 +635,14 @@ llvm::Value *CodeGenLLVM::Visit(const ir::Block *op) {
llvm::Value *CodeGenLLVM::Visit(const ir::PrimitiveNode *) { CINN_NOT_IMPLEMENTED return nullptr; }

llvm::Value *CodeGenLLVM::Visit(const ir::Call *op) {
LOG(INFO) << "function name=" << op->name;
if (op->name == runtime::intrinsic::debug_log_repr) {
return EmitCall_debug_info(op);
} else if (op->is_extern_call()) {
LOG(INFO) << op->name << " is extern call.";
auto emitter_id = ExternFuncID{backend_llvm_host, op->name.c_str()};
const auto &fn_name = ExternFunctionEmitterRegistry::Global().Lookup(emitter_id);
LOG(INFO) << fn_name;
if (!fn_name.empty()) {
ExternFunctionLLVMEmitter emitter(fn_name);
emitter.BindCodeGen(this);
Expand Down
10 changes: 10 additions & 0 deletions cinn/backends/llvm/llvm_intrin_rule.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -102,6 +102,16 @@ void RegisterCpuIntrinRule() {
}
});

ir::Registry::Register("lower_cpu_intrinsic_rsqrt", true).SetBody([](lang::Args args, lang::RetValue *rv) {
CHECK_GE(args.size(), 1U);
Expr arg0 = args[0];
ir::Call *node = arg0->as<ir::Call>();
CHECK(node);
CHECK(!node->read_args.empty());
Expr arg = node->read_args[0];
*rv = make_const(arg->type(), 1) / lang::Sqrt(arg);
});

ir::Registry::Register("lower_cpu_intrinsic_exp10", true).SetBody([](lang::Args args, lang::RetValue *rv) {
CHECK_GE(args.size(), 1U);
Expr arg0 = args[0];
Expand Down
188 changes: 171 additions & 17 deletions cinn/common/cas.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -141,6 +141,15 @@ bool IsDivisible(int64_t a, int64_t b) {
return a % b == 0;
}
bool IsDivisible(const Sum* a, int b);

// If int a Divisiable to any operands of product b
bool IsDivisible(int a, const Product* b) {
if (a < 0) return false;
for (auto& item : b->operands()) {
if (item.As<IntImm>() && item.As<IntImm>()->value > 0 && IsDivisible(a, item.As<IntImm>()->value)) return true;
}
return false;
}
bool IsDivisible(const Product* a, int b) {
for (auto& item : a->operands()) {
if (item.As<IntImm>() && IsDivisible(item.As<IntImm>()->value, b)) return true;
Expand Down Expand Up @@ -198,14 +207,28 @@ Expr Divide(const Product* a, int b) {
break;
}
}
// NOTE that a should be divisible by b.
CHECK(is_divisible) << "a should be divisible by b";
if (times != 1) {
args.push_back(make_const(a->type(), times));
}
for (int j = 0; j < a->operands().size(); j++) {
if (j == i) continue;
args.push_back(a->operand(j));
// Case is_divisible : a = 8x and b = 4 and a/b = 2x
// Case !is_divisible : a = 2x and b = 8 and a/b = x/4
if (is_divisible) {
// NOTE that a should be divisible by b.
if (times != 1) {
args.push_back(make_const(a->type(), times));
}
for (int j = 0; j < a->operands().size(); j++) {
if (j == i) continue;
args.push_back(a->operand(j));
}
return Product::Make(args);
} else {
for (i = 0; i < a->operands().size(); i++) {
auto* a_i = a->operand(i).As<IntImm>();
if (a_i && b % a_i->value == 0) {
b = b / a_i->value;
} else {
args.push_back(a->operand(i));
}
}
return FracOp::Make(Product::Make(args), Expr(b));
}
return Product::Make(args);
}
Expand Down Expand Up @@ -1117,16 +1140,39 @@ bool CasSimplifyMutator::GetMaxBound(Expr* lower_bound, Expr* upper_bound, Expr
}

bool CasSimplifyMutator::SimplifySpecificSumMod(Expr* result, Expr a, Expr b) {
// case1: (32+(-x))%33 = 32-x%33 (0<=x<=32)
// case2: (x-32))%33 = x%33 - 32%33 (0<=x<=32)
#ifdef CINN_WITH_CUDA
return false;
#else
// case1: (32+(-x))%33 = 32-x%33 (0<=x<=32)
// case2: (x-32))%33 = x%33 - 32%33 (0<=x<=32)
auto a_sum = a.As<Sum>();
auto b_i = b.As<IntImm>();
if (!a_sum || !b_i) {
return false;
}
// if 0 < b < 3, (3a+b) % 6 = (3a % 6) + (b % 6)
if (a_sum->operands().size() == 2) {
a_sum->operands()[0] = CasSimplify(a_sum->operands()[0], var_intervals);
auto sum_a_prod = a_sum->operands()[0].As<Product>();
auto sum_b_var = a_sum->operands()[1].As<_Var_>();
if (sum_a_prod && sum_b_var && var_intervals.count(sum_b_var->name)) {
auto sum_a_prod_b_int = sum_a_prod->operand(1).As<IntImm>();
if (sum_a_prod_b_int) std::swap(sum_a_prod->operand(0), sum_a_prod->operand(1));
auto sum_a_prod_a_int = sum_a_prod->operand(0).As<IntImm>();
auto& interval = var_intervals.at(sum_b_var->name);
int b_abs = std::abs(b_i->value);
int sum_prod_a_abs = std::abs(sum_a_prod_a_int->value);
if (sum_a_prod_a_int && (b_abs % sum_prod_a_abs == 0)) {
if (std::abs(interval.l) < sum_prod_a_abs && std::abs(interval.r) < sum_prod_a_abs) {
*result = CasSimplify(Sum::Make({CasSimplify(Mod::Make(a_sum->operands()[0], b), var_intervals),
CasSimplify(Mod::Make(a_sum->operands()[1], b), var_intervals)}),
var_intervals);
return true;
}
}
}
}
#ifdef CINN_WITH_CUDA
return false;
#else

int const_value = 0;
Expr lower_bound;
Expr upper_bound;
Expand Down Expand Up @@ -1216,9 +1262,26 @@ Expr CasSimplifyMutator::SimplifyMod(Expr u) {
// x % 1 = 0
if (b_i && b_i->value == 1) return make_const(b_i->type(), 0);

// 2x % 2 = 0
if (b_i && a_product && a_product->operand(0).As<IntImm>()) {
// 4x % 2 = 0
// 2x % 4 = 2 * (x % 2)
if (b_i && a_product && a_product->operand(0).As<IntImm>() && b_i->value > 0 &&
a_product->operand(0).As<IntImm>()->value > 0) {
if (a_product->operand(0).As<IntImm>()->value % b_i->value == 0) return make_const(a_product->type(), 0);
if (b_i->value % a_product->operand(0).As<IntImm>()->value == 0) {
int a_product_int = a_product->operand(0).As<IntImm>()->value;
int new_b = b_i->value / a_product_int;
return Product::Make({Expr(a_product_int), SimplifyMod(Mod::Make(a_product->operand(1), Expr(new_b)))});
}
}

if (b_i && a_product && a_product->operand(1).As<IntImm>() && b_i->value > 0 &&
a_product->operand(0).As<IntImm>()->value > 0) {
if (a_product->operand(1).As<IntImm>()->value % b_i->value == 0) return make_const(a_product->type(), 0);
if (b_i->value % a_product->operand(1).As<IntImm>()->value == 0) {
int a_product_int = a_product->operand(1).As<IntImm>()->value;
int new_b = b_i->value / a_product_int;
return Product::Make({Expr(a_product_int), SimplifyMod(Mod::Make(a_product->operand(0), Expr(new_b)))});
}
}

// 0 % x = 1, 1 % x = 1
Expand Down Expand Up @@ -1650,6 +1713,74 @@ Expr ConvertCinnToCAS(Expr expr) {
return copied;
}

/**
* @brief Given an expr, visit it. If there is an ir::Min and its operands are 1 constant value and 1 inconstant value,
* return the constant min value.
* For example, if a < min(5, b), then we get a < 5 and a < b. Using a < 5 to simplify the condition ensures
* correctness, though not sufficient.
*/
Expr ReplaceMinToConstant(Expr expr) {
Expr copied = optim::IRCopy(expr);
struct Mutator : public ir::IRMutator<ir::Expr*> {
void operator()(Expr* expr) { Visit(expr); }
void Visit(Expr* expr) { ir::IRMutator<>::Visit(expr, expr); }

private:
void Visit(const Min* op, Expr* expr) override {
auto a = op->a();
auto b = op->b();

Visit(&a);
Visit(&b);

auto min_a = op->a();
auto min_b = op->b();
if (min_a.is_constant() && !min_b.is_constant()) {
CHECK(min_a->type().is_integer());
*expr = optim::IRCopy(min_a);
} else if (min_b.is_constant() && !min_a.is_constant()) {
CHECK(min_b->type().is_integer());
*expr = optim::IRCopy(min_b);
}
}
};
Mutator()(&copied);
return copied;
}

/**
* @brief Given an expr, visit it. If there is an ir::Max and its operands are 1 constant value and 1 inconstant value,
* return the constant max value.
*/
Expr ReplaceMaxToConstant(Expr expr) {
Expr copied = optim::IRCopy(expr);
struct Mutator : public ir::IRMutator<ir::Expr*> {
void operator()(Expr* expr) { Visit(expr); }
void Visit(Expr* expr) { ir::IRMutator<>::Visit(expr, expr); }

private:
void Visit(const Max* op, Expr* expr) override {
auto a = op->a();
auto b = op->b();

Visit(&a);
Visit(&b);

auto max_a = op->a();
auto max_b = op->b();
if (max_a.is_constant() && !max_b.is_constant()) {
CHECK(max_a->type().is_integer());
*expr = optim::IRCopy(max_a);
} else if (max_b.is_constant() && !max_a.is_constant()) {
CHECK(max_b->type().is_integer());
*expr = optim::IRCopy(max_b);
}
}
};
Mutator()(&copied);
return copied;
}

Expr ConvertCasToCinn(Expr expr) {
Expr copied = optim::IRCopy(expr);

Expand Down Expand Up @@ -1800,7 +1931,7 @@ Expr DividePartially(Sum* a, int b) {
std::vector<Expr> external_sum_args, sum_args;

for (auto& item : a->operands()) {
if (item.As<Product>() && IsDivisible(item.As<Product>(), b)) {
if (item.As<Product>() && (IsDivisible(item.As<Product>(), b) || IsDivisible(b, item.As<Product>()))) {
external_sum_args.push_back(Divide(item.As<Product>(), b));
} else if (item.As<IntImm>() && IsDivisible(item.As<IntImm>()->value, b)) {
external_sum_args.push_back(make_const(item.type(), item.As<IntImm>()->value / b));
Expand Down Expand Up @@ -1930,12 +2061,35 @@ Expr CasSimplifyMutator::SimplifyFracOp(Expr expr) {
// disiviable
if (a_sum && IsDivisible(a_sum, bi->value)) return Divide(a_sum, bi->value);
if (a_product) {
if (IsDivisible(a_product, bi->value)) {
if (IsDivisible(a_product, bi->value) || IsDivisible(bi->value, a_product)) {
return Divide(a_product, bi->value);
} else {
return FracOp::Make(a, b);
}
}

// if 0 < b < 3, (3a+b) / 6 = (3a / 6) + (b / 6)
if (a_sum && a_sum->operands().size() == 2) {
a_sum->operands()[0] = CasSimplify(a_sum->operands()[0], var_intervals);
auto sum_a_prod = a_sum->operands()[0].As<Product>();
auto sum_b_var = a_sum->operands()[1].As<_Var_>();
if (sum_a_prod && sum_b_var && var_intervals.count(sum_b_var->name)) {
auto sum_a_prod_b_int = sum_a_prod->operand(1).As<IntImm>();
if (sum_a_prod_b_int) std::swap(sum_a_prod->operand(0), sum_a_prod->operand(1));
auto sum_a_prod_a_int = sum_a_prod->operand(0).As<IntImm>();
auto& interval = var_intervals.at(sum_b_var->name);
int b_abs = std::abs(bi->value);
int sum_prod_a_abs = std::abs(sum_a_prod_a_int->value);
if (sum_a_prod_a_int && (b_abs % sum_prod_a_abs == 0)) {
if (std::abs(interval.l) < sum_prod_a_abs && std::abs(interval.r) < sum_prod_a_abs) {
return CasSimplify(Sum::Make({CasSimplify(FracOp::Make(a_sum->operands()[0], b), var_intervals),
CasSimplify(FracOp::Make(a_sum->operands()[1], b), var_intervals)}),
var_intervals);
}
}
}
}

// not divisiable
/*
if (a_sum) {
Expand Down
38 changes: 35 additions & 3 deletions cinn/common/cas.h
100644 → 100755
View file
Original file line number Diff line number Diff line change
Expand Up @@ -21,10 +21,16 @@

#include "cinn/ir/ir.h"
#include "cinn/ir/ir_printer.h"
#include "cinn/optim/ir_simplify.h"

namespace cinn {
namespace common {

namespace detail {
Expr ReplaceMinToConstant(Expr expr);
Expr ReplaceMaxToConstant(Expr expr);
} // namespace detail

/**
* Interval of a _Var_.
*/
Expand All @@ -33,16 +39,42 @@ struct CasInterval {
CasInterval(T l, T r) : l(l), r(r) {
CHECK_LE(l, r) << "left shoud not be larger than right";
}
CasInterval(Expr e_l, Expr e_r) : e_l(e_l), e_r(e_r) {}

/**
* @brief When iterator's upper_bound is an ir::Min of a constant value and a inconstant value, choose the constant
* value. When iterator's lower_bound is an ir::Max of a constant value and a inconstant value, choose the constant
* value. E.g: expr_l = max(x, 1) and expr_r = min(y,5): max(x, 1) <= iterator_i <= min(y,5)
*
* the bounds will be simplified to e_l = 1 and e_r = 5:
* 1 <= iterator_i <= 5
*/
CasInterval(Expr expr_l, Expr expr_r) {
VLOG(2) << "CasInterval is : [" << expr_l << ", " << expr_r << "].";
expr_r = detail::ReplaceMinToConstant(expr_r);
expr_l = detail::ReplaceMaxToConstant(expr_l);
optim::Simplify(&expr_l);
optim::Simplify(&expr_r);
VLOG(2) << "After simplify, CasInterval is : [" << expr_l << ", " << expr_r << "].";

if (expr_l.is_constant() && expr_r.is_constant()) {
CHECK(expr_l->type().is_integer());
CHECK(expr_r->type().is_integer());
l = expr_l.as_int32();
r = expr_r.as_int32();
return;
}
e_l = expr_l;
e_r = expr_r;
}
int l, r;
// Note: not verify l <= r and (e_l, e_r) has higher priority than (l, r)
Expr e_l, e_r;

friend std::ostream& operator<<(std::ostream& os, const CasInterval& i) {
if (i.e_l.defined() && i.e_r.defined()) {
os << "Interval[" << i.e_l << ", " << i.e_r << "]";
os << "Expr e_l Interval[" << i.e_l << ", " << i.e_r << "]";
} else {
os << "Interval[" << i.l << ", " << i.r << "]";
os << "Int l Interval[" << i.l << ", " << i.r << "]";
}
return os;
}
Expand Down
3 changes: 3 additions & 0 deletions cinn/frontend/base_builder.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -37,6 +37,9 @@ class BaseBuilder {
// name of this builder
const std::string& name() { return name_; }

// the number of instructions
const size_t size() { return instrs_.size(); }

virtual ~BaseBuilder() {}

void AppendInstruction(const Instruction& instr) { instrs_.push_back(instr); }
Expand Down
Loading

0 comments on commit 3221e6a

Please sign in to comment.