[GitHub] [tvm] yzh119 commented on a change in pull request #8466: [CUDA] Initial support for dynamic shared memory

[GitBox <gi...@apache.org>]

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: commits-unsubscribe@tvm.apache.org

For queries about this service, please contact Infrastructure at:
users@infra.apache.org