yzh119 commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r669776990
##########
File path: src/runtime/thread_storage_scope.h
##########
@@ -42,21 +42,23 @@ enum class StorageRank {
kGlobal = 0,
/*! \brief shared memory among thread group */
kShared = 1,
+ /*! \brief dynamic shared memory among thread group */
+ kDynShared = 2,
/*!
* \brief reserved for warp memory.
* This is only used by programming model.
* There is no such memory usually in GPU.
* Instead, we can simulate it by registers and shuffle.
*/
- kWarp = 2,
+ kWarp = 3,
Review comment:
Would this break backward compatibility?
##########
File path: src/tir/transforms/split_host_device.cc
##########
@@ -89,6 +92,17 @@ class VarUseDefAnalysis : public StmtExprMutator {
Stmt VisitStmt_(const AllocateNode* op) final {
this->HandleDef(op->buffer_var.get());
+ auto storage_scope =
runtime::StorageScope::Create(GetPtrStorageScope(op->buffer_var));
+ if (storage_scope.rank == runtime::StorageRank::kDynShared) {
+ ICHECK_EQ(use_dyn_shmem_, false) << "Only one dynamic shared memory
allocation is allowed.";
Review comment:
Shall we support allocating multiple arrays from dynamically allocated
shared memory in the future?
```cuda
extern __shared__ float buffer[];
float* arr1 = (float*)buffer;
int* arr2 = (int*)&arr1[128];
...
```
--
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]
