Địa chỉ của bài viết này: https://www.cnblogs.com/wanger-sjtu/p/17501119.html.

PASS của VectorizeLoop là vector hóa vòng lặp For được đánh dấu là ForKind::kVectorized và thay thế các biến liên quan đến các câu lệnh trong vòng lặp For bằng Ramp để tạo điều kiện thuận lợi cho việc tạo các hướng dẫn vận hành vector hóa liên quan trong quá trình Codegen.

Chức năng nhập của VectorizeLoop PASS như sau. Tải xuống sẽ chỉ được kích hoạt khi Enable_vectorize=true được bật, nếu không VectorizeSkipper sẽ thay thế vòng lặp For của ForKind::kVectorized bằng vòng lặp bình thường.

                        
                          Truyền VectorizeLoop(bool enable_vectorize) { tự động pass_func = [=](PrimFunc f, IRModule m, PassContext ctx) { tự động* n = f.CopyOnWrite(); nếu (enable_vectorize) { n->body = LoopVectorizer()(std::move(n->body)); } nếu không { n->body = VectorizeSkipper()(std::move(n->body)); } trả về f; }; trả về CreatePrimFuncPass(pass_func, 0, "tir.VectorizeLoop", {}); }

                        
                      

Hãy lấy một vài ví dụ trong UT để giới thiệu cách triển khai mã nguồn.

vòng lặp vectorize

                        
                          dtype = "int64" n = te.var("n") ib = tvm.tir.ir_builder.create() A = ib.pointer("float32", name="A") với ib.for_range(0, n) là i: với ib.for_range(0, 4, kind="vectorize") là j: A[i*4+j] += tvm.tir.const(1, A.dtype) stmt = ib.get() assert isinstance(stmt.body, tvm.tir.For) mod = tvm.IRModule.from_expr(tvm.tir.PrimFunc([A, n], stmt)) stmt = tvm.tir.transform.VectorizeLoop()(mod)["main"].body

                        
                      

Đoạn mã trên thực hiện phép cộng vectơ. Vectơ A có độ dài 4n cộng +1 cho mỗi phần tử.

                        
                          # trước đối với (i, 0, n) { vectơ hóa (j, 0, 4) { A[((i*4) + j)] = (A[((i*4) + j)] + 1f) } } # sau đối với (i, 0, n) { A[ramp((i*4), 1, 4)] = (A[ramp((i*4), 1, 4)] + x4(1f)) }

                        
                      

Có thể thấy rằng sau khi vượt qua PASS của VectorizeLoop, vòng lặp bên trong bị loại bỏ và thay thế bằng lệnh vectơ Ramp. Lệnh này sẽ được thay thế bằng lệnh SIMD (neon, AVX, v.v.) trong CPU.

quá trình ĐẠT

PASS được vector hóa được xử lý trong LoopVectorizer, xử lý phần vòng lặp For.

                        
                          lớp LoopVectorizer: public StmtMutator { public: Stmt VisitStmt_(const ForNode* op) final { if (op->kind == ForKind::kVectorized) { ICHECK(is_zero(op->min)); auto* extend_as_int = op->extent.as(); if (!extent_as_int || extend_as_int->value < 1) { LOG(FATAL) << "Không thể vector hóa vòng lặp với phạm vi " << op->extent; } return Vectorizer(op->loop_var, static_cast(extent_as_int->value))(op->body); } else { return StmtMutator::VisitStmt_(op); } } };


                        
                      

Khi gặp một nút cần được vector hóa, trước tiên hãy ghi lại biến và phạm vi vòng lặp, những biến này sẽ được sử dụng sau này khi thay thế các hoạt động Tải và Lưu trữ tương ứng bằng Ramp. Sau đó, chúng ta đến phần Vectorizer, duyệt qua phần thân vòng lặp For và sửa đổi stmt tương ứng.

                        
                          Vectorizer(Var var, int var_lanes) : var_(var), var_lanes_(var_lanes) { đoạn đường nối_ = Đường dốc(0, 1, var_lanes); }

                        
                      

Các PrimExpr và Stmt khác nhau bị quá tải trong Vectorizer. Tôi sẽ không giới thiệu từng cái một ở đây mà sẽ sử dụng phép tính cộng vectơ ở trên để giới thiệu các hàm và quy trình được sử dụng.

Trước tiên, chúng ta hãy xem logic tính toán trong vòng lặp For trong sch ở trên:

                        
                           A[((i*4) + j)] = (A[((i*4) + j)] + 1f)

                        
                      

Bởi vì trong TVM, biểu thức của Stmt có thể được coi là ngôn ngữ DSL và AST được duyệt theo chiến lược theo chiều sâu khi được truy cập. Ở đây, quy trình tính toán ở trên được biểu thị đơn giản dưới dạng cây cú pháp AST và sau đó là các lệnh gọi. trong quá trình này được phân tích mỗi chức năng được xử lý như thế nào.

Như có thể thấy từ sơ đồ AST ở trên, đối với sch trên, BufferStoreNode, Add Mul, BufferLoadNode, v.v. được truy cập theo trình tự. Ở đây chúng tôi sẽ giới thiệu quy trình vector hóa dựa trên việc xử lý các Nút này.

Cái gọi là quá trình vector hóa là ánh xạ phép toán vòng lặp vô hướng được đánh dấu kVectorized thành một phép toán được vector hóa. Đối với ví dụ trên, tất cả các quyền truy cập vào j đều được ánh xạ tới RampNode, để quá trình xử lý tiếp theo có thể tạo ra các hướng dẫn tương ứng một cách chính xác.

Nút lưu trữ đệm

Có ba phần trong BufferStoreNode:

• bộ đệm——bộ đệm để ghi vào
• giá trị - giá trị hoặc biểu thức được viết
• chỉ số——tọa độ được ghi vào bộ đệm
  Mục đích ở đây là sửa đổi giá trị Và chỉ số nội dung trong.
  vì chỉ số , được thực hiện ở đây. cuối cùng đã vượt qua Trợ lý bản đồ Lần lượt ghé thăm chỉ số sự biểu lộ.

                        
                          auto fmutate = [this](const PrimExpr& index) { return this->VisitExpr(index); }; Mảng chỉ số = op->indices.Map(fmutate);

                        
                      

Đối với giá trị, nó được duyệt trực tiếp.

                        
                          Giá trị PrimExpr = this->VisitExpr(op->giá trị);

                        
                      

Thêm nút

Cả AddNode và SubNode đều sẽ chuyển đến chức năng mẫu AddSubVec. Hàm này trước tiên sẽ duyệt qua các biểu thức bên trái và bên phải, .

                        
                          PrimExpr a = this->VisitExpr(op->a); PrimExpr b = this->VisitExpr(op->b); nếu (a. giống như(op->a) && b. giống như(op->b)) { trả về GetRef(op); } else { int lanes = std::max(a.dtype().lanes(), b.dtype().lanes()); nếu (lanes != 1) { const RampNode* b_ramp = b.as(); const RampNode* a_ramp = a.as(); nếu (a.dtype().lanes() == 1 && b_ramp) { trả về Ramp(fcompute(a, b_ramp->base), fcompute(make_zero(b_ramp->stride.dtype()), b_ramp->stride), b_ramp->lanes); } nếu (b.dtype().lanes() == 1 && a_ramp) { trả về Ramp(fcompute(a_ramp->base, b), a_ramp->stride, a_ramp->lanes); } } trả về fcompute(BroadcastTo(a, lanes), BroadcastTo(b, lanes));

                        
                      

Nếu không có thay đổi sau khi duyệt qua, nó sẽ được trả lại trực tiếp. Và điều chúng ta cần tính toán ở đây là.

                        
                          ((i*4) + j)

                        
                      

j là tọa độ cần được vector hóa. i*4 không thay đổi. Sau khi truyền tải, a không thay đổi và b trở thành T.Ramp(0, 1, 4). Tại thời điểm này, làn đường = 4, nó sẽ chuyển đến nhánh if đầu tiên và RampNode mới được xây dựng sẽ được trả về.

                        
                           T.Ramp(i * 4, 1, 4)

                        
                      

Các nhánh khác cũng tương tự. Ví dụ:

                        
                          A[i * 4 + j] + T.float32(1) // --- sau --- A[i * 4:i * 4 + 4] T.float32(1)

                        
                      

Ở đây, a và b sẽ được phát dưới dạng vectơ rồi tính toán.

Mã số

Phán quyết của VarNode ở đây tương đối đơn giản. Nếu biến cần được vector hóa khớp với nhau thì RampNode được xây dựng trong hàm tạo sẽ được trả về, nếu không nó sẽ được trả về. Các hoạt động khác được bỏ qua bây giờ.

                        
                          Var var = GetRef(op); if (var.same_as(var_)) { return đoạn đường nối_; } // ... else { return std::move(var); }

                        
                      

Nút đa năng

                        
                          PrimExpr a = this->VisitExpr(op->a); PrimExpr b = this->VisitExpr(op->b); nếu (a. giống như(op->a) && b. giống như(op->b)) { trả về GetRef(op); } else { int lanes = std::max(a.dtype().lanes(), b.dtype().lanes()); nếu (lanes != 1) { const RampNode* b_ramp = b.as(); const RampNode* a_ramp = a.as(); nếu (a_ramp && b.dtype().lanes() == 1 && analyzer_.CanProve(b > 0)) { trả về Ramp(a_ramp->base * b, a_ramp->stride * b, a_ramp->lanes); } nếu (b_ramp && a.dtype().lanes() == 1 && analyzer_.CanProve(a > 0)) { trả về Ramp(b_ramp->base * a, b_ramp->stride * a, b_ramp->lanes); } } trả về Mul(BroadcastTo(a, lanes), BroadcastTo(b, lanes)); } trả về BinaryVec(op);

                        
                      

Logic xử lý ở đây về cơ bản giống như Add. Nó chỉ khác một chút khi tính toán RampNode.

Cuối cùng, bài viết về đọc mã nguồn TVM PASS—VectorizeLoop kết thúc tại đây. Nếu bạn muốn biết thêm về cách đọc mã nguồn TVM PASS—VectorizeLoop, vui lòng tìm kiếm bài viết CFSDN hoặc tiếp tục duyệt qua các bài viết liên quan. blog tương lai! .