nverke commented on code in PR #13719:
URL: https://github.com/apache/tvm/pull/13719#discussion_r1070169824
##########
python/tvm/tir/tensor_intrin/hexagon.py:
##########
@@ -20,6 +20,47 @@
from .. import TensorIntrin
+def generate_dma_load_intrin(
+ size: int,
+ dtype: str,
+):
+ """Generator of dma_load intrins"""
+
+ @T.prim_func
+ def dma_load_desc(a: T.handle, c: T.handle) -> None:
+ A = T.match_buffer(a, (size), dtype, offset_factor=1, scope="global")
+ C = T.match_buffer(c, (size), dtype, offset_factor=1,
scope="global.vtcm")
+ with T.block("root"):
+ T.reads(A[0:size])
+ T.writes(C[0:size])
+ for i in T.serial(size):
+ with T.block("load"):
+ vii = T.axis.remap("S", [i])
+ C[vii] = A[vii]
+
+ @T.prim_func
+ def dma_load_impl(a: T.handle, c: T.handle) -> None:
+ A = T.match_buffer(a, (size), dtype, offset_factor=1, scope="global")
+ C = T.match_buffer(c, (size), dtype, offset_factor=1,
scope="global.vtcm")
+ with T.block("root"):
+ T.reads(A[0:size])
+ T.writes(C[0:size])
+ T.evaluate(
+ T.tvm_call_packed(
+ "device_api.hexagon.dma_copy",
+ 0,
Review Comment:
✅
##########
python/tvm/tir/tensor_intrin/hexagon.py:
##########
@@ -20,6 +20,47 @@
from .. import TensorIntrin
+def generate_dma_load_intrin(
+ size: int,
+ dtype: str,
+):
+ """Generator of dma_load intrins"""
+
+ @T.prim_func
+ def dma_load_desc(a: T.handle, c: T.handle) -> None:
Review Comment:
✅
##########
python/tvm/tir/tensor_intrin/hexagon.py:
##########
@@ -20,6 +20,47 @@
from .. import TensorIntrin
+def generate_dma_load_intrin(
+ size: int,
+ dtype: str,
+):
+ """Generator of dma_load intrins"""
+
+ @T.prim_func
+ def dma_load_desc(a: T.handle, c: T.handle) -> None:
+ A = T.match_buffer(a, (size), dtype, offset_factor=1, scope="global")
+ C = T.match_buffer(c, (size), dtype, offset_factor=1,
scope="global.vtcm")
+ with T.block("root"):
+ T.reads(A[0:size])
+ T.writes(C[0:size])
+ for i in T.serial(size):
+ with T.block("load"):
+ vii = T.axis.remap("S", [i])
+ C[vii] = A[vii]
+
+ @T.prim_func
+ def dma_load_impl(a: T.handle, c: T.handle) -> None:
Review Comment:
✅
##########
python/tvm/tir/tensor_intrin/hexagon.py:
##########
@@ -20,6 +20,47 @@
from .. import TensorIntrin
+def generate_dma_load_intrin(
+ size: int,
+ dtype: str,
+):
+ """Generator of dma_load intrins"""
+
+ @T.prim_func
+ def dma_load_desc(a: T.handle, c: T.handle) -> None:
+ A = T.match_buffer(a, (size), dtype, offset_factor=1, scope="global")
+ C = T.match_buffer(c, (size), dtype, offset_factor=1,
scope="global.vtcm")
+ with T.block("root"):
+ T.reads(A[0:size])
+ T.writes(C[0:size])
+ for i in T.serial(size):
+ with T.block("load"):
+ vii = T.axis.remap("S", [i])
+ C[vii] = A[vii]
+
+ @T.prim_func
+ def dma_load_impl(a: T.handle, c: T.handle) -> None:
+ A = T.match_buffer(a, (size), dtype, offset_factor=1, scope="global")
+ C = T.match_buffer(c, (size), dtype, offset_factor=1,
scope="global.vtcm")
+ with T.block("root"):
+ T.reads(A[0:size])
+ T.writes(C[0:size])
+ T.evaluate(
+ T.tvm_call_packed(
+ "device_api.hexagon.dma_copy",
+ 0,
+ T.address_of(C[0], dtype="handle"),
+ T.address_of(A[0], dtype="handle"),
+ size,
+ 0,
Review Comment:
Going to just add a comment about this for now. These intrins don't
currently have any planned use so figure if we find one we can add increased
functionality.
##########
python/tvm/tir/tensor_intrin/hexagon.py:
##########
@@ -20,6 +20,47 @@
from .. import TensorIntrin
+def generate_dma_load_intrin(
+ size: int,
+ dtype: str,
+):
+ """Generator of dma_load intrins"""
+
+ @T.prim_func
+ def dma_load_desc(a: T.handle, c: T.handle) -> None:
+ A = T.match_buffer(a, (size), dtype, offset_factor=1, scope="global")
+ C = T.match_buffer(c, (size), dtype, offset_factor=1,
scope="global.vtcm")
+ with T.block("root"):
+ T.reads(A[0:size])
+ T.writes(C[0:size])
+ for i in T.serial(size):
+ with T.block("load"):
+ vii = T.axis.remap("S", [i])
+ C[vii] = A[vii]
+
+ @T.prim_func
+ def dma_load_impl(a: T.handle, c: T.handle) -> None:
+ A = T.match_buffer(a, (size), dtype, offset_factor=1, scope="global")
+ C = T.match_buffer(c, (size), dtype, offset_factor=1,
scope="global.vtcm")
+ with T.block("root"):
+ T.reads(A[0:size])
+ T.writes(C[0:size])
+ T.evaluate(
+ T.tvm_call_packed(
+ "device_api.hexagon.dma_copy",
+ 0,
+ T.address_of(C[0], dtype="handle"),
+ T.address_of(A[0], dtype="handle"),
+ size,
+ 0,
+ dtype="int32",
+ )
+ )
+ T.evaluate(T.tvm_call_packed("device_api.hexagon.dma_wait", 0, 0,
dtype="int32"))
Review Comment:
✅
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
