You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by GitBox <gi...@apache.org> on 2021/11/03 03:21:07 UTC
[GitHub] [tvm] SYangDong opened a new pull request #9436: add is_entry_func tag for device function
SYangDong opened a new pull request #9436:
URL: https://github.com/apache/tvm/pull/9436
Add KIsEntryFunc tag for the device function after being split. Normally, the device function after being split should be a entry function. With this attribute, the global fucniton or device funtion can be generated by checking the KIsEntryFunc tag. That will be useful when generaing cross function calls.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tqchen edited a comment on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
tqchen edited a comment on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-960058753
I agree we should avoid entry function concept in kernel functions(and only use it for main). Most cases we should just be explicit about what the function is. If there are motivation for other use cases "global" vs "device", let us use another tag
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] comaniac commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
comaniac commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-958655466
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tqchen commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
tqchen commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-960058753
I agree we should avoid entry function concept in kernel functions. Most cases we should just be explicit about what the function is. If there are motivation for other use cases "global" vs "device", let us use another tag
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] comaniac commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
comaniac commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-958655466
I don't think this is an entry function by definition. Device functions can be more flexible and are not limited to one entry. For example, we could have multiple device functions and all of them can be invoked directly by the host. In this case, none of the device functions should be only entry function.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Hzfengsy commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
Hzfengsy commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-958635367
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] vinx13 commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
vinx13 commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-959818875
I think what you need is a tag to differentiate `global` and `device` functions. Even if in a single operator, it is possible that multiple kernel launches are needed. My concern is that reusing `kIsEntryFunc` may break the assumption that entry func is unique
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Hzfengsy commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
Hzfengsy commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-958635367
IIUC, the device function is not the entry function.
Take `cuda` backend as an example, we start the program on the host (`llvm` program) and than call `cuda` kernels. So the host function is the entry function.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] SYangDong commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
SYangDong commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-958642319
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] SYangDong commented on pull request #9436: Add is_global_func tag to differentiate global and device function
Posted by GitBox <gi...@apache.org>.
SYangDong commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-963013599
@vinx13 The test case has been added. Please have a look.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] SYangDong edited a comment on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
SYangDong edited a comment on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-958676448
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] SYangDong edited a comment on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
SYangDong edited a comment on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-958676448
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] SYangDong commented on pull request #9436: Add is_global_func tag to differentiate global and device function
Posted by GitBox <gi...@apache.org>.
SYangDong commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-960424292
@vinx13 @tqchen Thank you for providing more details about the definition of kIsEntryFunc. I try to use another tag. Please have a look if you have the bandwidth.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] comaniac commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
comaniac commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-958655466
I don't think this is an entry function by definition. Device functions can be more flexible and are not limited to one entry. For example, we could have multiple device functions and all of them can be invoked directly by the host. In this case, none of the device functions should be only entry function.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Hzfengsy commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
Hzfengsy commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-958635367
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] SYangDong edited a comment on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
SYangDong edited a comment on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-958676448
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tqchen edited a comment on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
tqchen edited a comment on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-960058753
I agree we should avoid entry function concept in kernel functions(and only use it for main). Most cases we should just be explicit about what the function is. If there are motivation for other use cases "global" vs "device", let us use another tag
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] SYangDong commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
SYangDong commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-958642319
@Hzfengsy Thank you for a quick reply. I got what you mean. It depends how you define the entry function. The device function after being split also can be called entry function for this device module. With this tag, we can easily check __device__ or __global__.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] SYangDong edited a comment on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
SYangDong edited a comment on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-958676448
If we consider only one single operator, the purpose of this modification is to mark the device function after being split as entry function(map to the ``__global__`` function in cuda code). The other PrimFuncs in the IRModule are ``__device_`` functions in code code.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] SYangDong commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
SYangDong commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-958642319
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Hzfengsy commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
Hzfengsy commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-958635367
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] vinx13 commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
vinx13 commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-959818875
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tqchen edited a comment on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
tqchen edited a comment on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-960058753
I agree we should avoid entry function concept in kernel functions(and only use it for main). Most cases we should just be explicit about what the function is. If there are motivation for other use cases "global" vs "device", let us use another tag
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] SYangDong edited a comment on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
SYangDong edited a comment on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-958676448
If we consider only one single operator, the purpose of this modification is to mark the device function after being split as entry function(map to the global function in cuda code). The other PrimFuncs in the IRModule are ``__device_`` functions in code code.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] SYangDong commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
SYangDong commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-958676448
If we consider only one single operator, the purpose of this modification is to mark the device function after being split as entry function(map to the global function in cuda code). The other PrimFuncs in the IRModule are` `__device_`` functions in code code.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Hzfengsy commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
Hzfengsy commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-958644956
cc @junrushao1994 @vinx13
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] vinx13 commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
vinx13 commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-959818875
I think what you need is a tag to differentiate `global` and `device` functions. Even if in a single operator, it is possible that multiple kernel launches are needed. My concern is that reusing `kIsEntryFunc` may break the assumption that entry func is unique
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tqchen commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
tqchen commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-960058753
I agree we should avoid entry function concept in kernel functions. Most cases we should just be explicit about what the function is. If there are motivation for other use cases "global" vs "device", let us use another tag
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] SYangDong commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
SYangDong commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-958668046
@comaniac Thank you for your explaination. IMHO, entry function is a relative concept. Relative to host funciton，the device function is not entry function，definitely. But when having __device__ function and __global__ fucntion, the global function is the entry function of the device function. Maybe I misuse the tag entry_function. Take the code below for example:
```A = te.var('A')
B = te.var('B')
callee = tir.PrimFunc([A, B], tir.Evaluate(tir.Add(A, B)))
callee = callee.with_attr('global_symbol', 'callee')
main = tir.PrimFunc([A, B], tir.Evaluate(tir.Call('int32', callee, [A, B])))
# OR main = tir.PrimFunc([A, B], tir.Evaluate(tir.Call('int32', GlobalVar('callee'), [A, B])))
main = main.with_attr("global_symbol", "main")
main = main.with_attr("tir.noalias", True)
mod = tvm.IRModule({"main": main, "callee": callee})
tvm.build(mod)
```
So if I add a tag of the main function, tvm can generate the "__global__ void" adn "__device__ void" correspondingly.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] vinx13 merged pull request #9436: Add is_global_func tag to differentiate global and device function
Posted by GitBox <gi...@apache.org>.
vinx13 merged pull request #9436:
URL: https://github.com/apache/tvm/pull/9436
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] vinx13 commented on pull request #9436: Add is_global_func tag to differentiate global and device function
Posted by GitBox <gi...@apache.org>.
vinx13 commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-960424782
Could you add a test case for SplitHostDevice?
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tqchen commented on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
tqchen commented on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-960058753
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tqchen edited a comment on pull request #9436: Add is_entry_func tag for device function
Posted by GitBox <gi...@apache.org>.
tqchen edited a comment on pull request #9436:
URL: https://github.com/apache/tvm/pull/9436#issuecomment-960058753
I agree we should avoid entry function concept in kernel functions(and only use it for main). Most cases we should just be explicit about what the function is. If there are motivation for other use cases "global" vs "device", let us use another tag
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org