tvm tutorial (1.2)

最新推荐文章于 2022-03-18 15:28:26 发布
最新推荐文章于 2022-03-18 15:28:26 发布
阅读量207 收藏
点赞数
分类专栏： LLVM/TVM/编译理论 文章标签： python
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/tengfei461807914/article/details/122822236
版权

cache writing

紧接着上一篇实现了array packing对矩阵乘法进行优化的策略，接下来记录cache写回优化（block writing throung caching)

cache的写回说头比较多，有两种情况，一种叫write back，另一种是write through，详细内容可以参考下面链接：
https://blog.csdn.net/xingzhe22222/article/details/81988101

矩阵C的写回操作也是跳跃式的，不利于cache的更新操作。因此这里将计算结果保存在一个临时空间当中，并以顺序存储。

tvm代码如下，代码中暂时去掉了向量化和并行化的策略，生成的IR更清晰一些：

 packedB = te.compute((N / bn, K, bn), lambda bigN, k, littleN: B[k, bigN * bn + littleN], name="packedB")
 C = te.compute(
     (M, N),
     lambda x, y: te.sum(A[x, k] * packedB[y // bn, k, tvm.tir.indexmod(y, bn)], axis=k),
     name="C",
 )
 s = te.create_schedule(C.op)

 # Allocate write cache
 CC = s.cache_write(C, "global")

 xo, yo, xi, yi = s[C].tile(C.op.axis[0], C.op.axis[1], bn, bn)

 # Write cache is computed at yo
 s[CC].compute_at(s[C], yo)

 # New inner axes
 xc, yc = s[CC].op.axis

 (k,) = s[CC].op.reduce_axis
 ko, ki = s[CC].split(k, factor=4)
 s[CC].reorder(ko, xc, ki, yc)

 print(tvm.lower(s, [A, B, C], simple_mode=True))

生成的IR如下：

primfn(A_1: handle, B_1: handle, C_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {A: Buffer(A_2: Pointer(float32), float32, [1024, 1024], []),
             C: Buffer(C_2: Pointer(float32), float32, [1024, 1024], []),
             B: Buffer(B_2: Pointer(float32), float32, [1024, 1024], [])}
  buffer_map = {A_1: A, B_1: B, C_1: C} {
  allocate(packedB: Pointer(global float32), float32, [1048576]), storage_scope = global;
  allocate(C.global: Pointer(global float32), float32, [1024]), storage_scope = global {
    for (bigN: int32, 0, 32) {
      for (k: int32, 0, 1024) {
        for (littleN: int32, 0, 32) {
          packedB[(((bigN*32768) + (k*32)) + littleN)] = (float32*)B_2[(((k*1024) + (bigN*32)) + littleN)]
        }
      }
    }
    for (x.outer: int32, 0, 32) {
      for (y.outer: int32, 0, 32) {
        for (x.c.init: int32, 0, 32) {
          for (y.c.init: int32, 0, 32) {
            C.global[((x.c.init*32) + y.c.init)] = 0f32
          }
        }
        for (k.outer: int32, 0, 256) {
          for (x.c: int32, 0, 32) {
            for (k.inner: int32, 0, 4) {
              for (y.c: int32, 0, 32) {
                C.global[((x.c*32) + y.c)] = ((float32*)C.global[((x.c*32) + y.c)] + ((float32*)A_2[((((x.outer*32768) + (x.c*1024)) + (k.outer*4)) + k.inner)]*(float32*)packedB[((((y.outer*32768) + (k.outer*128)) + (k.inner*32)) + y.c)]))
              }
            }
          }
        }
        for (x.inner: int32, 0, 32) {
          for (y.inner: int32, 0, 32) {
            C_2[((((x.outer*32768) + (x.inner*1024)) + (y.outer*32)) + y.inner)] = (float32*)C.global[((x.inner*32) + y.inner)]
          }
        }
      }
    }
  }
}

s[CC].cache_write 这条语句在文档中的表述如下：

Create a cache write of original tensor, before storing into tensor.
This will mutate the body of the tensor. A new cache stage will created before feed into the tensor.
This function can be used to support data layout transformation. If there is a split/fuse/reorder on the data parallel axis of tensor before cache_write is called. The intermediate cache stores the data in the layout as the iteration order of leave axis. The data will be transformed back to the original layout in the original tensor. User can further call compute_inline to inline the original layout and keep the data stored in the transformed layout.

文档中的第一句话就表达了含义，"在存储到tensor之前，创建一个原始张量的写入cache“
可以理解成“开辟了”一个cache的空间专门用于存储输出结果（CPU上的cache是用户不可见的）

（简单翻译一下）
这个函数可以用来支持数据layout的变换，如果在cache写入操作被执行之前，在数据可以并行化的axis上存在split/fuse/reorder这些操作，intermediate cache将存储这些数据以结束对应axis迭代的数据摆放顺序。数据也会转换会原始tensor中的数据layout形式。用户可以调用compute_inline来实现对原始数据layout的inline，并保持数据以变换后的layout。

总结一下，这个函数可以变换数据摆放形式，比如NCHW可以变成NHWC这样。
如果有split/fuse/reorder这样的操作，数据在cache当中会按照数据的连续迭代数据摆放，使用完后会转换回原有的layout.

这里对比一下生成的IR，将执行这条语句下面的所有schedule全部注释掉，得到的IR如下：

primfn(A_1: handle, B_1: handle, C_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {C: Buffer(C_2: Pointer(float32), float32, [1024, 1024], []),
             A: Buffer(A_2: Pointer(float32), float32, [1024, 1024], []),
             B: Buffer(B_2: Pointer(float32), float32, [1024, 1024], [])}
  buffer_map = {A_1: A, B_1: B, C_1: C} {
  allocate(packedB: Pointer(global float32), float32, [1048576]), storage_scope = global;
  allocate(C.global: Pointer(global float32), float32, [1048576]), storage_scope = global {
    for (bigN: int32, 0, 32) {
      for (k: int32, 0, 1024) {
        for (littleN: int32, 0, 32) {
          packedB[(((bigN*32768) + (k*32)) + littleN)] = (float32*)B_2[(((k*1024) + (bigN*32)) + littleN)]
        }
      }
    }
    for (x.c: int32, 0, 1024) {
      for (y.c: int32, 0, 1024) {
        C.global[((x.c*1024) + y.c)] = 0f32
        for (k_1: int32, 0, 1024) {
          C.global[((x.c*1024) + y.c)] = ((float32*)C.global[((x.c*1024) + y.c)] + ((float32*)A_2[((x.c*1024) + k_1)]*(float32*)packedB[(((floordiv(y.c, 32)*32768) + (k_1*32)) + floormod(y.c, 32))]))
        }
      }
    }
    for (x: int32, 0, 1024) {
      for (y: int32, 0, 1024) {
        C_2[((x*1024) + y)] = (float32*)C.global[((x*1024) + y)]
      }
    }
  }
}

去掉这条语句后得到的IR为：

primfn(A_1: handle, B_1: handle, C_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {A: Buffer(A_2: Pointer(float32), float32, [1024, 1024], []),
             C: Buffer(C_2: Pointer(float32), float32, [1024, 1024], []),
             B: Buffer(B_2: Pointer(float32), float32, [1024, 1024], [])}
  buffer_map = {A_1: A, B_1: B, C_1: C} {
  allocate(packedB: Pointer(global float32), float32, [1048576]), storage_scope = global {
    for (bigN: int32, 0, 32) {
      for (k: int32, 0, 1024) {
        for (littleN: int32, 0, 32) {
          packedB[(((bigN*32768) + (k*32)) + littleN)] = (float32*)B_2[(((k*1024) + (bigN*32)) + littleN)]
        }
      }
    }
    for (x: int32, 0, 1024) {
      for (y: int32, 0, 1024) {
        C_2[((x*1024) + y)] = 0f32
        for (k_1: int32, 0, 1024) {
          C_2[((x*1024) + y)] = ((float32*)C_2[((x*1024) + y)] + ((float32*)A_2[((x*1024) + k_1)]*(float32*)packedB[(((floordiv(y, 32)*32768) + (k_1*32)) + floormod(y, 32))]))
        }
      }
    }
  }
}

去掉不相干的IR，得到如下两个对比，最下面的是使用s[CC].cache_write

    for (x: int32, 0, 1024) {
      for (y: int32, 0, 1024) {
        C_2[((x*1024) + y)] = 0f32
        for (k_1: int32, 0, 1024) {
          C_2[((x*1024) + y)] = ((float32*)C_2[((x*1024) + y)] + ((float32*)A_2[((x*1024) + k_1)]*(float32*)packedB[(((floordiv(y, 32)*32768) + (k_1*32)) + floormod(y, 32))]))
        }
      }
    }
  }
}

  allocate(C.global: Pointer(global float32), float32, [1048576]), storage_scope = global {
    for (x.c: int32, 0, 1024) {
      for (y.c: int32, 0, 1024) {
        C.global[((x.c*1024) + y.c)] = 0f32
        for (k_1: int32, 0, 1024) {
          C.global[((x.c*1024) + y.c)] = ((float32*)C.global[((x.c*1024) + y.c)] + ((float32*)A_2[((x.c*1024) + k_1)]*(float32*)packedB[(((floordiv(y.c, 32)*32768) + (k_1*32)) + floormod(y.c, 32))]))
        }
      }
    }
    for (x: int32, 0, 1024) {
      for (y: int32, 0, 1024) {
        C_2[((x*1024) + y)] = (float32*)C.global[((x*1024) + y)]
      }
    }
  }
}

很明显哈，创建了一个临时变量C.global用来计算，计算完成后再写回到原来的tensor C_2当中

再看一个语句compute_at函数
官方文档给出的解释如下：

Attach the stage at parent’s scope

绑定到作用于的某一个循环轴上，对比一下增加与不增加这条语句得到的IR的区别（这里把tile部分加上）
代码如下：

 s = te.create_schedule(C.op)
 CC = s.cache_write(C, "global")
 xo, yo, xi, yi = s[C].tile(C.op.axis[0], C.op.axis[1], bn, bn)
 s[CC].compute_at(s[C], yo)

生成的IR如下：

primfn(A_1: handle, B_1: handle, C_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {C: Buffer(C_2: Pointer(float32), float32, [1024, 1024], []),
             A: Buffer(A_2: Pointer(float32), float32, [1024, 1024], []),
             B: Buffer(B_2: Pointer(float32), float32, [1024, 1024], [])}
  buffer_map = {A_1: A, B_1: B, C_1: C} {
  allocate(packedB: Pointer(global float32), float32, [1048576]), storage_scope = global;
  allocate(C.global: Pointer(global float32), float32, [1024]), storage_scope = global {
    for (bigN: int32, 0, 32) {
      for (k: int32, 0, 1024) {
        for (littleN: int32, 0, 32) {
          packedB[(((bigN*32768) + (k*32)) + littleN)] = (float32*)B_2[(((k*1024) + (bigN*32)) + littleN)]
        }
      }
    }
    for (x.outer: int32, 0, 32) {
      for (y.outer: int32, 0, 32) {
        for (x.c: int32, 0, 32) {
          for (y.c: int32, 0, 32) {
            C.global[((x.c*32) + y.c)] = 0f32
            for (k_1: int32, 0, 1024) {
              C.global[((x.c*32) + y.c)] = ((float32*)C.global[((x.c*32) + y.c)] + ((float32*)A_2[(((x.outer*32768) + (x.c*1024)) + k_1)]*(float32*)packedB[(((y.outer*32768) + (k_1*32)) + y.c)]))
            }
          }
        }
        for (x.inner: int32, 0, 32) {
          for (y.inner: int32, 0, 32) {
            C_2[((((x.outer*32768) + (x.inner*1024)) + (y.outer*32)) + y.inner)] = (float32*)C.global[((x.inner*32) + y.inner)]
          }
        }
      }
    }
  }
}

当仅使用tile情况下的代码为：

 s = te.create_schedule(C.op)
 CC = s.cache_write(C, "global")
 xo, yo, xi, yi = s[C].tile(C.op.axis[0], C.op.axis[1], bn, bn)

生成的IR如下：

primfn(A_1: handle, B_1: handle, C_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {A: Buffer(A_2: Pointer(float32), float32, [1024, 1024], []),
             C: Buffer(C_2: Pointer(float32), float32, [1024, 1024], []),
             B: Buffer(B_2: Pointer(float32), float32, [1024, 1024], [])}
  buffer_map = {A_1: A, B_1: B, C_1: C} {
  allocate(packedB: Pointer(global float32), float32, [1048576]), storage_scope = global;
  allocate(C.global: Pointer(global float32), float32, [1048576]), storage_scope = global {
    for (bigN: int32, 0, 32) {
      for (k: int32, 0, 1024) {
        for (littleN: int32, 0, 32) {
          packedB[(((bigN*32768) + (k*32)) + littleN)] = (float32*)B_2[(((k*1024) + (bigN*32)) + littleN)]
        }
      }
    }
    for (x.c: int32, 0, 1024) {
      for (y.c: int32, 0, 1024) {
        C.global[((x.c*1024) + y.c)] = 0f32
        for (k_1: int32, 0, 1024) {
          C.global[((x.c*1024) + y.c)] = ((float32*)C.global[((x.c*1024) + y.c)] + ((float32*)A_2[((x.c*1024) + k_1)]*(float32*)packedB[(((floordiv(y.c, 32)*32768) + (k_1*32)) + floormod(y.c, 32))]))
        }
      }
    }
    for (x.outer: int32, 0, 32) {
      for (y.outer: int32, 0, 32) {
        for (x.inner: int32, 0, 32) {
          for (y.inner: int32, 0, 32) {
            C_2[((((x.outer*32768) + (x.inner*1024)) + (y.outer*32)) + y.inner)] = (float32*)C.global[((((x.outer*32768) + (x.inner*1024)) + (y.outer*32)) + y.inner)]
          }
        }
      }
    }
  }
}

对比最后将C.global写回到C_2部分可以发现，调用s[CC].compute_at(s[C], yo)函数后，会将写回的逻辑嵌套再C的计算逻辑中执行，并在会嵌套再yo的下面

最后补上向量化和循环展开以及并行化部分的逻辑，TVM的代码如下：

s = te.create_schedule(C.op)

# Allocate write cache
CC = s.cache_write(C, "global")

xo, yo, xi, yi = s[C].tile(C.op.axis[0], C.op.axis[1], bn, bn)

# Write cache is computed at yo
s[CC].compute_at(s[C], yo)

# New inner axes
xc, yc = s[CC].op.axis

(k,) = s[CC].op.reduce_axis
ko, ki = s[CC].split(k, factor=4)
s[CC].reorder(ko, xc, ki, yc)
s[CC].unroll(ki)
s[CC].vectorize(yc)

x, y, z = s[packedB].op.axis
s[packedB].vectorize(z)
s[packedB].parallel(x)

得到的IR如下：

primfn(A_1: handle, B_1: handle, C_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {C: Buffer(C_2: Pointer(float32), float32, [1024, 1024], []),
             A: Buffer(A_2: Pointer(float32), float32, [1024, 1024], []),
             B: Buffer(B_2: Pointer(float32), float32, [1024, 1024], [])}
  buffer_map = {A_1: A, B_1: B, C_1: C} {
  allocate(packedB: Pointer(global float32x32), float32x32, [32768]), storage_scope = global;
  allocate(C.global: Pointer(global float32), float32, [1024]), storage_scope = global {
    for (bigN: int32, 0, 32) "parallel" {
      for (k: int32, 0, 1024) {
        packedB[ramp(((bigN*32768) + (k*32)), 1, 32)] = (float32x32*)B_2[ramp(((k*1024) + (bigN*32)), 1, 32)]
      }
    }
    for (x.outer: int32, 0, 32) {
      for (y.outer: int32, 0, 32) {
        for (x.c.init: int32, 0, 32) {
          C.global[ramp((x.c.init*32), 1, 32)] = broadcast(0f32, 32)
        }
        for (k.outer: int32, 0, 256) {
          for (x.c: int32, 0, 32) {
            C.global[ramp((x.c*32), 1, 32)] = ((float32x32*)C.global[ramp((x.c*32), 1, 32)] + (broadcast((float32*)A_2[(((x.outer*32768) + (x.c*1024)) + (k.outer*4))], 32)*(float32x32*)packedB[ramp(((y.outer*32768) + (k.outer*128)), 1, 32)]))
            C.global[ramp((x.c*32), 1, 32)] = ((float32x32*)C.global[ramp((x.c*32), 1, 32)] + (broadcast((float32*)A_2[((((x.outer*32768) + (x.c*1024)) + (k.outer*4)) + 1)], 32)*(float32x32*)packedB[ramp((((y.outer*32768) + (k.outer*128)) + 32), 1, 32)]))
            C.global[ramp((x.c*32), 1, 32)] = ((float32x32*)C.global[ramp((x.c*32), 1, 32)] + (broadcast((float32*)A_2[((((x.outer*32768) + (x.c*1024)) + (k.outer*4)) + 2)], 32)*(float32x32*)packedB[ramp((((y.outer*32768) + (k.outer*128)) + 64), 1, 32)]))
            C.global[ramp((x.c*32), 1, 32)] = ((float32x32*)C.global[ramp((x.c*32), 1, 32)] + (broadcast((float32*)A_2[((((x.outer*32768) + (x.c*1024)) + (k.outer*4)) + 3)], 32)*(float32x32*)packedB[ramp((((y.outer*32768) + (k.outer*128)) + 96), 1, 32)]))
          }
        }
        for (x.inner: int32, 0, 32) {
          for (y.inner: int32, 0, 32) {
            C_2[((((x.outer*32768) + (x.inner*1024)) + (y.outer*32)) + y.inner)] = (float32*)C.global[((x.inner*32) + y.inner)]
          }
        }
      }
    }
  }
}

并行化

并行化的操作再之前的优化中已经使用过，代码如下：

s = te.create_schedule(C.op)

# Allocate write cache
CC = s.cache_write(C, "global")

xo, yo, xi, yi = s[C].tile(C.op.axis[0], C.op.axis[1], bn, bn)

# Write cache is computed at yo
s[CC].compute_at(s[C], yo)

# 并行化
s[C].parallel(xo)

# New inner axes
xc, yc = s[CC].op.axis

(k,) = s[CC].op.reduce_axis
ko, ki = s[CC].split(k, factor=4)
s[CC].reorder(ko, xc, ki, yc)
s[CC].unroll(ki)
s[CC].vectorize(yc)

x, y, z = s[packedB].op.axis
s[packedB].vectorize(z)
s[packedB].parallel(x)

代码中只对矩阵C的schedule上增加了一个对M外侧循环的并行化处理的指令

得到的IR如下：

primfn(A_1: handle, B_1: handle, C_1: handle) -> ()
  attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True}
  buffers = {C: Buffer(C_2: Pointer(float32), float32, [1024, 1024], []),
             A: Buffer(A_2: Pointer(float32), float32, [1024, 1024], []),
             B: Buffer(B_2: Pointer(float32), float32, [1024, 1024], [])}
  buffer_map = {A_1: A, B_1: B, C_1: C} {
  allocate(packedB: Pointer(global float32x32), float32x32, [32768]), storage_scope = global {
    for (bigN: int32, 0, 32) "parallel" {
      for (k: int32, 0, 1024) {
        packedB[ramp(((bigN*32768) + (k*32)), 1, 32)] = (float32x32*)B_2[ramp(((k*1024) + (bigN*32)), 1, 32)]
      }
    }
    for (x.outer: int32, 0, 32) "parallel" {
      allocate(C.global: Pointer(global float32), float32, [1024]), storage_scope = global;
      for (y.outer: int32, 0, 32) {
        for (x.c.init: int32, 0, 32) {
          C.global[ramp((x.c.init*32), 1, 32)] = broadcast(0f32, 32)
        }
        for (k.outer: int32, 0, 256) {
          for (x.c: int32, 0, 32) {
            C.global[ramp((x.c*32), 1, 32)] = ((float32x32*)C.global[ramp((x.c*32), 1, 32)] + (broadcast((float32*)A_2[(((x.outer*32768) + (x.c*1024)) + (k.outer*4))], 32)*(float32x32*)packedB[ramp(((y.outer*32768) + (k.outer*128)), 1, 32)]))
            C.global[ramp((x.c*32), 1, 32)] = ((float32x32*)C.global[ramp((x.c*32), 1, 32)] + (broadcast((float32*)A_2[((((x.outer*32768) + (x.c*1024)) + (k.outer*4)) + 1)], 32)*(float32x32*)packedB[ramp((((y.outer*32768) + (k.outer*128)) + 32), 1, 32)]))
            C.global[ramp((x.c*32), 1, 32)] = ((float32x32*)C.global[ramp((x.c*32), 1, 32)] + (broadcast((float32*)A_2[((((x.outer*32768) + (x.c*1024)) + (k.outer*4)) + 2)], 32)*(float32x32*)packedB[ramp((((y.outer*32768) + (k.outer*128)) + 64), 1, 32)]))
            C.global[ramp((x.c*32), 1, 32)] = ((float32x32*)C.global[ramp((x.c*32), 1, 32)] + (broadcast((float32*)A_2[((((x.outer*32768) + (x.c*1024)) + (k.outer*4)) + 3)], 32)*(float32x32*)packedB[ramp((((y.outer*32768) + (k.outer*128)) + 96), 1, 32)]))
          }
        }
        for (x.inner: int32, 0, 32) {
          for (y.inner: int32, 0, 32) {
            C_2[((((x.outer*32768) + (x.inner*1024)) + (y.outer*32)) + y.inner)] = (float32*)C.global[((x.inner*32) + y.inner)]
          }
        }
      }
    }
  }
}

总结

te的schedule提供了一系列用于计算优化的操作，如多核并行、循环展开以及循环分割等等，用户可以针对不同的数据规模和数据形状提供手动的优化。

样例代码中提供了te模块用于优化的使用步骤：

  1. 使用一系列operation来描述对数据的计算流程，即使用lambda表达式实现计算逻辑。
  2. 使用schedule primitives描述如何计算，即进行上面的优化过程。
  3. 编译到不同的平台上进行codegen
  4. 如果需要，可以把生成的IR以及机器码存储起来
红鲤鱼遇绿鲤鱼
关注
  • 0
    点赞
  • 踩
  • 0
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
专栏目录
【tvm官网教程】张量表达与调度
B10030607的博客
04-21 1925
【tvm官网教程】张量表达与调度目的1. 调度原语1.1 te常用接口1.2 stage常用成员函数欢迎使用Markdown编辑器新的改变功能快捷键合理的创建标题，有助于目录的生成如何改变文本的样式插入链接与图片如何插入一段漂亮的代码片生成一个适合你的列表创建一个表格设定内容居中、居左、居右SmartyPants创建一个自定义列表如何创建一个注脚注释也是必不可少的KaTeX数学公式新的甘特图功能，丰富你的文章UML 图表FLowchart流程图导出与导入导出导入 目的 编译器是下一代大厂之争的关键，曾经有过
TVM Compiler中文教程：TVM调度原语（Schedule Primitives）
Mars_WH的专栏
06-10 8018
文章目录TVM调度原语（Schedule Primitives）分裂split平铺tile TVM调度原语（Schedule Primitives） TVM是用于高效内核代码构建的版本领域专用语言（Domain-Specialed-Language，DSL） 。 这篇教程，我们将展示通过TVM提供的各种原语怎么去调度计算。 from __future__ import absolute_impor...
tvm tutorial (1.1)
@fei
02-08 1408
接着上次 tensor_expr_get_started.py 没记录完的部分继续 剩下的代码如下，给出的样例是Nvidia GPU设备端的代码 run_cuda = False if run_cuda: # Change this target to the correct backend for you gpu. For example: cuda (NVIDIA GPUs), # rocm (Radeon GPUS), OpenCL (opencl). tgt_gpu = tv
tvm tutorials 记录（1）
@fei
02-03 1105
第一篇文章记录了relay_quick_start.py文件中的内容，主要是展示了如何使用relay构建网络结构图，以及调用什么接口可以实现对网络进行编译优化并部署。 接下来看一下tensor_expr_get_started.py这个文件，文件900多行，注释占了绝大部分，提供这个demo的作者人很实在。 作者上来来了一段解释，原文如下： In this tutorial we will turn our attention to how TVM works with Tensor Expression
TVM之Schedule
sayandroid的专栏
07-07 4042
TVM之Schedule1. schedule2. stage3. methods for stage3.1 split3.2 tile3.3 fuse3.4 reorder3.5 bind3.6 compute_at3.7 compute_inline3.8 compute_root 最近在阅读TVM官网的文档，在阅读的过程中将keywords以及相应理解记录了下来，有理解不对的地方，欢迎批评指...
tvm_3rdparty
07-05
《TVM中的3rdparty模块详解》 在深度学习领域，TVM是一个高效且高度优化的开源跨平台运行时系统，用于部署机器学习模型。它以其灵活的编译器架构和广泛的硬件支持而受到赞誉。在TVM的源码结构中，`3rdparty`目录...
TVM jvmpkg 编译完成包
04-14
版本是V8.0以后的版本，源码于2022年4月中旬下载并编译通过，包含tvm中整个完整jvm的内容以及全部生成内容，方便因为编译环境备受折磨的小伙伴直接下载使用。其中包括jvm/core/target 下编译android端需要的jar包库...
tvm.tar TVM源码压缩包
09-24
TVM完整源码，自己下载，若网速不好，可能会有部分链接缺失，导致编译出错，血泪教训我知道！！！！
tvm.rar_tvm_轨道TVM全拼
09-19
标题中的“tvm.rar_tvm_轨道TVM全拼”暗示了我们正在讨论的是轨道交通行业的自动售票机（Ticket Vending Machine，简称TVM）的技术细节，特别是与单片机控制输入输出相关的控制逻辑。在轨道交通系统中，TVM是乘客...
TVM share-lib build so库包
04-14
《TVM Share-Lib：构建SO库包的深度解析》 在现代的机器学习领域，TVM（TVM：The Virtual Machine for Deep Learning）扮演着至关重要的角色。它是一个开源的深度学习编译器框架，旨在优化和部署各种硬件平台上的...
TVM User Tutorial -- Introduction to TOPI
zachariah2000的博客
03-06 394
Author: Ehsan M. Kermani 这是 TVM Operator Inventory (TOPI) 的介绍性教程。 TOPI 提供了 numpy 风格的通用操作和调度，其抽象程度高于 TVM。 在本教程中，我们将了解 TOPI 如何使我们免于在 TVM 中编写模板代码。 from __future__ import absolute_import, print_function import tvm import tvm.testing from tvm import te from t
TVM（一）：简介与安装
热门推荐
u014448054的博客
09-25 1万+
简介 TVM是一个用于深度学习系统的编译器堆栈。 它旨在缩小以生产力为中心的深度学习框架与以性能和效率为中心的硬件后端之间的差距。TVM与深度学习框架合作，为不同的后端提供端到端编译。 换句话说，TVM就是一种将深度学习工作负载部署到硬件的端到端IR（中间表示）堆栈。可以表述为一种把深度学习模型分发到各种硬件设备上的、端到端的解决方案。 GitHub：https://github.com/dmlc...
TVM源码分析——总体安排
weixin_45803667的博客
12-14 2787
TVM源码分析（一）简介TVM源码目录结构后续博客安排 简介 本系列博客的主要目的是博主在打算学习TVM时，发现对TVM的源码介绍内容非常少，因此本系列博客主要以学习TVM框架为前提下，来对TVM的源码内容进行解读，从而也能帮助更多的人来了解TVM。 TVM源码目录结构 3rdparty：主要是TVM使用到的第三方库 apps：包含了TVM的扩展项目，并有使用TVM的示例 cmake：cmake...
初识 TVM
小菜鸟的AI之路
11-15 6661
如有图像或公式显示错误，可以访问我的个人博客：https://www.wanglichun.tech/2019/11/15/tvm/ 笔者也是最近偶然的机会才开始接触TVM，使用过后发现，经过auto-tuning后的TVM模型在速度是竟然超过了TensorRT,并且笔者使用的是MXNet框架，TVM对MXNet绝对的友好，对于Pytorch等模型，可以使用ONNX，操作也一样简单，使用起来基本类...
TVM tutorials代码学习-from_darknet.py
Mars_WH的专栏
10-05 989
TVM tutorials代码学习-from_darknet.py准备环境详解小结 准备环境 要运行tvm源码下tutorials/frontend/from_darknet.py代码需要额外安装cffi和cv2包 pip install cffi pip install opencv-python 详解 #导入包 import numpy as np #计算库 import matpl...
使用TVM优化深度学习GPU算子：深度卷积实例
yiran103的专栏
03-01 1万+
以下内容翻译自：Optimize Deep Learning GPU Operators with TVM: A Depthwise Convolution Example 高效的深度学习算子是深度学习系统的核心。通常这些算子很难优化，并且需要高性能计算专家的努力。TVM，端到端张量IR/DSL堆栈，使得这项任务更容易。 这个博客教你如何在TVM的帮助下编写高性能GPU运算核心。我们使用深度卷...
tvm 入门（二）
qq_39790992的博客
05-24 2638
代码是一个tvm入门的例子。 以向量相加为例。 使用tvm的流程是： 1.描述串行的向量相加是怎么做的 2.描述并行的时候，怎么对计算单元做划分 3.编译目标函数。本文所示代码可以看到用tvm生成的cuda代码。 4.把编译生成的内容保存成文件，加载该文件做高性能计算。（可选） """ Get Started with TVM ==================== **Aut...
如何在 CPU 上优化 GEMM
吴建明wujianming_110117
10-05 378
如何在 CPU 上优化 GEMM (TL;DR) TVM 提供抽象接口，允许用户分别描述算法和算法的实施组织（所谓的调度）。通常，在高性能调度中编写算法，会破坏算法的可读性和模块化。尝试各种看似有前途的调度也很耗时。在 TVM 的帮助下，可以有效地尝试这些调度，提高性能。 将演示如何使用 TVM 优化矩阵乘法，通过简单地添加 18 行额外代码，实现比基线快 200 倍。 在 CPU 上执行的密集计算应用程序，有两个重要的优化： 提高内存访问的缓存命中率。复杂的数值计算和热点内存访问，都可以通过高缓存命中率
How to Guides -- How to optimize GEMM on CPU
zachariah2000的博客
03-18 275
Author: Jian Weng, Ruofei Yu TVM 提供抽象接口，允许用户分别描述算法和算法的实现组织（所谓的调度）。 通常，以高性能调度编写算法会破坏算法的可读性和模块化。 此外，尝试各种看似有希望的时间表非常耗时。 在 TVM 的帮助下，我们可以有效地尝试这些时间表以提高性能。 在本教程中，我们将演示如何使用 TVM 优化方阵乘法，并通过简单地添加 18 行额外的代码实现比基线快 200 倍。 在 CPU 上执行的密集计算应用程序有两个重要的优化： 提高内存访问的缓存命中率。 高缓存命中
tvm ubuntu
最新发布
09-10
安装与配置的步骤如下： 1. 首先，你需要下载并解压LLVM包（clang llvm-10.0.0-x86_64-linux-gnu-ubuntu-18.04.tar.xz），该包包含了所需的编译工具。 2. 解压后，你可以将文件夹重命名为clang-llvm-10.0.0，以方便后续的操作。 3. 接下来，你需要添加环境变量。可以在终端中执行以下命令将LLVM路径添加到环境变量中： ```shell export LLVM_PATH=${HOME}/tools/clang-llvm-10.0.0 ``` 4. 然后，你需要安装TVM。在打开Python解释器后，执行以下命令： ```python import tvm tvm.__version__ ``` 这将显示当前安装的TVM版本。 5. 最后，你需要克隆TVM的GitHub仓库，并执行以下命令： ```shell git clone --recursive https://github.com/apache/incubator-tvm tvm ``` 这将下载TVM的源代码并将其存储到名为tvm的文件夹中。 通过按照上述步骤进行操作，你将能够成功在Ubuntu上安装和配置TVM。
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值