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/07/14 02:28:43 UTC
[GitHub] [tvm] masahi opened a new pull request #8466: [CUDA] Initial support for dynamic shared memory
masahi opened a new pull request #8466:
URL: https://github.com/apache/tvm/pull/8466
Only one allocation is allowed for now.
--
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 a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
tqchen commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r672230286
##########
File path: src/runtime/file_utils.cc
##########
@@ -64,12 +66,14 @@ void FunctionInfo::Save(dmlc::Stream* writer) const {
writer->Write(name);
writer->Write(arg_types);
writer->Write(thread_axis_tags);
+ writer->Write(use_dyn_shared_memory);
}
bool FunctionInfo::Load(dmlc::Stream* reader) {
if (!reader->Read(&name)) return false;
if (!reader->Read(&arg_types)) return false;
if (!reader->Read(&thread_axis_tags)) return false;
+ if (!reader->Read(&use_dyn_shared_memory)) return false;
Review comment:
I think. let us think about another way to do backward compact then.
Since binary backward compatibility is more important to us here.
Here is one solution that also helps future backward compatibility.
Let us rename `thread_axis_tags => launch_param_tags`, and use `dyn_shared_memory` as a tag for the launching parameter of dynamic shared memory. We can then change the `ThreadAxisConfig => LaunchParamConfig`.
##########
File path: src/runtime/file_utils.cc
##########
@@ -64,12 +66,14 @@ void FunctionInfo::Save(dmlc::Stream* writer) const {
writer->Write(name);
writer->Write(arg_types);
writer->Write(thread_axis_tags);
+ writer->Write(use_dyn_shared_memory);
}
bool FunctionInfo::Load(dmlc::Stream* reader) {
if (!reader->Read(&name)) return false;
if (!reader->Read(&arg_types)) return false;
if (!reader->Read(&thread_axis_tags)) return false;
+ if (!reader->Read(&use_dyn_shared_memory)) return false;
Review comment:
Let us try the ^ change and do a test to make sure we can load back old cuda binaries
--
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] masahi commented on a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r672217797
##########
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:
I realized that we must support this usage, since it is required, among others, for multiplying fp16 or int8 matrices and accumulate into fp32 or int32, for example (tensor core).
I'm thinking that we need to expose pointer cast in TIR. I dind't find it in the current code:
```
extern __shared__ char buf[SIZE];
half* A_ptr = (half*)&buf[0];
half* B_ptr = (half*)&buf[128];
float32* C_ptr = (float32*)&buf[256];
```
I didn't get what you meant by "translate to an intrinsic" @tqchen. Can we hide pointer cast and the load /stores that follow it behind intrinsics? For allocation and merging, I'm thinking about working exclusively with byte arrays and make use of `storage_rewrite` as is.
--
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] masahi commented on a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r672217797
##########
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:
I realized that we must support this usage, since it is required, among others, for multiplying fp16 or int8 matrices and accumulate into fp32 or int32.
I'm thinking that we need to expose pointer cast in TIR. I dind't find it in the current code:
```
extern __shared__ char buf[SIZE];
half* A_ptr = (half*)&buf[0];
half* B_ptr = (half*)&buf[128];
float32* C_ptr = (float32*)&buf[256];
```
I didn't get what you meant by "translate to an intrinsic" @tqchen. Can we hide pointer cast and the load /stores that follow it behind intrinsics? For allocation and merging, I'm thinking about working exclusively with byte arrays and make use of `storage_rewrite` as is.
##########
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:
I realized that we must support this usage, since it is required, among others, for multiplying fp16 or int8 matrices and accumulate into fp32 or int32, for example (tensor core).
I'm thinking that we need to expose pointer cast in TIR. I dind't find it in the current code:
```
extern __shared__ char buf[SIZE];
half* A_ptr = (half*)&buf[0];
half* B_ptr = (half*)&buf[128];
float32* C_ptr = (float32*)&buf[256];
```
I didn't get what you meant by "translate to an intrinsic" @tqchen. Can we hide pointer cast and the load /stores that follow it behind intrinsics? For allocation and merging, I'm thinking about working exclusively with byte arrays and make use of `storage_rewrite` as is.
##########
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:
I realized that we must support this usage, since it is required, among others, for multiplying fp16 or int8 matrices and accumulate into fp32 or int32, for example (tensor core).
I'm thinking that we need to expose pointer cast in TIR, and somehow extend `Load` and `Store` node to be able to read/write with arbitrary pointers. I want to be able to express something like below :
```
extern __shared__ char buf[SIZE];
half* A_ptr = (half*)&buf[0];
half* B_ptr = (half*)&buf[128];
float32* C_ptr = (float32*)&buf[256];
...
C_ptr[0] = A_ptr[0] * B_ptr[0];
```
I didn't get what you meant by "translate to an intrinsic" @tqchen. Can we hide pointer cast and the load /stores that follow it behind intrinsics? For allocation and merging, I'm thinking about working exclusively with byte arrays and make use of `storage_rewrite` as is.
##########
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:
One possible solution is to rewrite the above program to
```
extern __shared__ char buf[SIZE];
*((floa32*)&buf[256]) = ((half*)&buf[0])[0] * ((half*)&buf[128])[0];
```
And introduce two intrinsics, `reinterpret_load` / `reinterpret_store`, which can be used to express the snippet above by
```
extern __shared__ char buf[SIZE];
reinterpret_store(&buf[256], "float32", 0,
reinterpret_load(&buf[0], "half", 0) * reinterpret_load(&buf[128], "half", 0));
```
##########
File path: src/runtime/file_utils.cc
##########
@@ -64,12 +66,14 @@ void FunctionInfo::Save(dmlc::Stream* writer) const {
writer->Write(name);
writer->Write(arg_types);
writer->Write(thread_axis_tags);
+ writer->Write(use_dyn_shared_memory);
}
bool FunctionInfo::Load(dmlc::Stream* reader) {
if (!reader->Read(&name)) return false;
if (!reader->Read(&arg_types)) return false;
if (!reader->Read(&thread_axis_tags)) return false;
+ if (!reader->Read(&use_dyn_shared_memory)) return false;
Review comment:
If we do the rename `thread_axis_tags => launch_param_tags`, the JSON reader would break backward compat at `helper.DeclareField("launch_param_tags", &launch_param_tags);` .
Or shall we do the rename but keep the old attribute name for JSON reader/writer? e.g.
`helper.DeclareField("thread_axis_tags", &launch_param_tags);`.
--
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] masahi commented on pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
masahi commented on pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#issuecomment-884664455
Thanks @tqchen @vinx13 @yzh119 @Hzfengsy
--
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 #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
tqchen commented on pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#issuecomment-881405667
Thanks @masahi I left some comments
--
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 #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
tqchen commented on pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#issuecomment-883405197
@masahi this is done by the code owner mechanism automatically by github. As long as there is one codeowner approval i think we are fine here.
--
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] masahi commented on pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
masahi commented on pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#issuecomment-883387968
Sorry apparently I made a bunch of review requests to code owners, but that was not intended.
--
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] masahi commented on a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r671689944
##########
File path: src/runtime/thread_storage_scope.h
##########
@@ -96,6 +98,8 @@ struct StorageScope {
return "global" + tag;
case StorageRank::kShared:
return "shared" + tag;
+ case StorageRank::kDynShared:
Review comment:
I changed dynamic shared memory to be a tagged memory, for `kShared`, but I had to add a check `e->scope.tag != ".dyn"` in a few places in `lower_device_storage_access_info.cc` and `storage_rewrite.cc` to workaround error/segfault.
--
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] masahi commented on pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
masahi commented on pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#issuecomment-883376561
@tqchen Backward compat issue has been addressed.
--
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 a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
tqchen commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r671194376
##########
File path: include/tvm/tir/function.h
##########
@@ -240,7 +240,7 @@ namespace attr {
*
* Call(f,
* [arg1, arg2, ..., arg_n,
- * work_size_1, work_size_2, ... work_size_m])
+ * work_size_1, work_size_2, ... work_size_m, dyn_shmem_size])
Review comment:
Let us introduce a special meta-data to indicate that dynamic shared memory is used. This is to make sure the calling convention is backward compatible when dyn shared memory is not provided.
```
constexpr const char* kDeviceUseDynSharedMemory = "tir.device_use_dyn_shared_memory";
```
##########
File path: src/runtime/thread_storage_scope.h
##########
@@ -223,6 +232,7 @@ class ThreadAxisConfig {
w.work_size[arg_index_map_[i]] = size;
}
}
+ w.dyn_shmem_size = static_cast<size_t>(x.values[base_ + arg_index_map_.size()].v_int64);
Review comment:
We need to consider backward compatibility for kernels where dynamic shared memory is not used.
Update https://github.com/apache/tvm/blob/main/src/runtime/meta_data.h#L103 to add a `use_dyn_shared_memory`(default to false) field, update json reader to use DeclareOptionalField, to optionally read this field in. Only decode the value when the flag is true.
##########
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:
We can do a lowering stage that translate allocate dynamic memory to an intrinsic that directly get the dynamic memory with an offset(where multiple allocation can be merged with proper alignment). This will help us get around some limitation imposed by storage rewrite. Split host device might be the right place for this.
--
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 a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
tqchen commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r672230575
##########
File path: src/runtime/file_utils.cc
##########
@@ -64,12 +66,14 @@ void FunctionInfo::Save(dmlc::Stream* writer) const {
writer->Write(name);
writer->Write(arg_types);
writer->Write(thread_axis_tags);
+ writer->Write(use_dyn_shared_memory);
}
bool FunctionInfo::Load(dmlc::Stream* reader) {
if (!reader->Read(&name)) return false;
if (!reader->Read(&arg_types)) return false;
if (!reader->Read(&thread_axis_tags)) return false;
+ if (!reader->Read(&use_dyn_shared_memory)) return false;
Review comment:
Let us try the ^ change and do a test to make sure we can load back old cuda binaries
--
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] masahi commented on a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r672235542
##########
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:
One possible solution is to rewrite the above program to
```
extern __shared__ char buf[SIZE];
*((floa32*)&buf[256]) = ((half*)&buf[0])[0] * ((half*)&buf[128])[0];
```
And introduce two intrinsics, `reinterpret_load` / `reinterpret_store`, which can be used to express the snippet above by
```
extern __shared__ char buf[SIZE];
reinterpret_store(&buf[256], "float32", 0,
reinterpret_load(&buf[0], "half", 0) * reinterpret_load(&buf[128], "half", 0));
```
--
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] masahi commented on a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r669976240
##########
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:
Should be no problem, as long as nobody depends on the integer value.
All tests passed, so I assume it is ok.
--
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] masahi commented on a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r673123145
##########
File path: src/runtime/thread_storage_scope.h
##########
@@ -194,13 +198,17 @@ struct ThreadWorkLoad {
inline size_t grid_dim(size_t i) const { return work_size[i]; }
};
/*! \brief Thread axis configuration */
-class ThreadAxisConfig {
+class LaunchParamConfig {
public:
- void Init(size_t base, const std::vector<std::string>& thread_axis_tags) {
+ void Init(size_t base, const std::vector<std::string>& launch_param_tags) {
base_ = base;
std::vector<bool> filled(6, false);
- for (size_t i = 0; i < thread_axis_tags.size(); ++i) {
- const std::string& tag = thread_axis_tags[i];
+ for (size_t i = 0; i < launch_param_tags.size(); ++i) {
+ const std::string& tag = launch_param_tags[i];
+ if (tag == kUseDynamicSharedMemoryTag) {
Review comment:
Added an assert
--
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] masahi commented on a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r669982717
##########
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:
I'm planning to support allocating multiple dynamic shmem at TIR level, like:
```
A_shared = ib.allocate(dtype, A.shape, scope="dyn.shared")
B_shared = ib.allocate(dtype, B.shape, scope="dyn.shared")
```
Under the hood, multiple allocations will be merged into a single allocation via `storage_rewrite.cc` pass. But I haven't thought about merging multiple allocations with hetero-dtype like your example.
Maybe for dynamic shmem, we just allocate a byte array and use `reinterpret` cast for splitting, assuming it is possible with TIR? Probably `storage_rewrite.cc` does not work for such scenario. Thought? @tqchen @junrushao1994 @vinx13
Also @yzh119 do you have a use case for such hetero-dtype dynamic shared memory?
--
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 a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
tqchen commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r673108940
##########
File path: src/runtime/file_utils.cc
##########
@@ -64,12 +66,14 @@ void FunctionInfo::Save(dmlc::Stream* writer) const {
writer->Write(name);
writer->Write(arg_types);
writer->Write(thread_axis_tags);
+ writer->Write(use_dyn_shared_memory);
}
bool FunctionInfo::Load(dmlc::Stream* reader) {
if (!reader->Read(&name)) return false;
if (!reader->Read(&arg_types)) return false;
if (!reader->Read(&thread_axis_tags)) return false;
+ if (!reader->Read(&use_dyn_shared_memory)) return false;
Review comment:
We can remove the thread_axis_tags later
--
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] masahi commented on a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r672217797
##########
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:
I realized that we must support this usage, since it is required, among others, for multiplying fp16 or int8 matrices and accumulate into fp32 or int32.
I'm thinking that we need to expose pointer cast in TIR. I dind't find it in the current code:
```
extern __shared__ char buf[SIZE];
half* A_ptr = (half*)&buf[0];
half* B_ptr = (half*)&buf[128];
float32* C_ptr = (float32*)&buf[256];
```
I didn't get what you meant by "translate to an intrinsic" @tqchen. Can we hide pointer cast and the load /stores that follow it behind intrinsics? For allocation and merging, I'm thinking about working exclusively with byte arrays and make use of `storage_rewrite` as is.
##########
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:
I realized that we must support this usage, since it is required, among others, for multiplying fp16 or int8 matrices and accumulate into fp32 or int32, for example (tensor core).
I'm thinking that we need to expose pointer cast in TIR. I dind't find it in the current code:
```
extern __shared__ char buf[SIZE];
half* A_ptr = (half*)&buf[0];
half* B_ptr = (half*)&buf[128];
float32* C_ptr = (float32*)&buf[256];
```
I didn't get what you meant by "translate to an intrinsic" @tqchen. Can we hide pointer cast and the load /stores that follow it behind intrinsics? For allocation and merging, I'm thinking about working exclusively with byte arrays and make use of `storage_rewrite` as is.
##########
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:
I realized that we must support this usage, since it is required, among others, for multiplying fp16 or int8 matrices and accumulate into fp32 or int32, for example (tensor core).
I'm thinking that we need to expose pointer cast in TIR, and somehow extend `Load` and `Store` node to be able to read/write with arbitrary pointers. I want to be able to express something like below :
```
extern __shared__ char buf[SIZE];
half* A_ptr = (half*)&buf[0];
half* B_ptr = (half*)&buf[128];
float32* C_ptr = (float32*)&buf[256];
...
C_ptr[0] = A_ptr[0] * B_ptr[0];
```
I didn't get what you meant by "translate to an intrinsic" @tqchen. Can we hide pointer cast and the load /stores that follow it behind intrinsics? For allocation and merging, I'm thinking about working exclusively with byte arrays and make use of `storage_rewrite` as is.
##########
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:
One possible solution is to rewrite the above program to
```
extern __shared__ char buf[SIZE];
*((floa32*)&buf[256]) = ((half*)&buf[0])[0] * ((half*)&buf[128])[0];
```
And introduce two intrinsics, `reinterpret_load` / `reinterpret_store`, which can be used to express the snippet above by
```
extern __shared__ char buf[SIZE];
reinterpret_store(&buf[256], "float32", 0,
reinterpret_load(&buf[0], "half", 0) * reinterpret_load(&buf[128], "half", 0));
```
--
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] masahi commented on a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r672217797
##########
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:
I realized that we must support this usage, since it is required, among others, for multiplying fp16 or int8 matrices and accumulate into fp32 or int32.
I'm thinking that we need to expose pointer cast in TIR. I dind't find it in the current code:
```
extern __shared__ char buf[SIZE];
half* A_ptr = (half*)&buf[0];
half* B_ptr = (half*)&buf[128];
float32* C_ptr = (float32*)&buf[256];
```
I didn't get what you meant by "translate to an intrinsic" @tqchen. Can we hide pointer cast and the load /stores that follow it behind intrinsics? For allocation and merging, I'm thinking about working exclusively with byte arrays and make use of `storage_rewrite` as is.
--
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 a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
tqchen commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r672230286
##########
File path: src/runtime/file_utils.cc
##########
@@ -64,12 +66,14 @@ void FunctionInfo::Save(dmlc::Stream* writer) const {
writer->Write(name);
writer->Write(arg_types);
writer->Write(thread_axis_tags);
+ writer->Write(use_dyn_shared_memory);
}
bool FunctionInfo::Load(dmlc::Stream* reader) {
if (!reader->Read(&name)) return false;
if (!reader->Read(&arg_types)) return false;
if (!reader->Read(&thread_axis_tags)) return false;
+ if (!reader->Read(&use_dyn_shared_memory)) return false;
Review comment:
I think. let us think about another way to do backward compact then.
Since binary backward compatibility is more important to us here.
Here is one solution that also helps future backward compatibility.
Let us rename `thread_axis_tags => launch_param_tags`, and use `dyn_shared_memory` as a tag for the launching parameter of dynamic shared memory. We can then change the `ThreadAxisConfig => LaunchParamConfig`.
--
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 a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
tqchen commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r673111931
##########
File path: src/runtime/thread_storage_scope.h
##########
@@ -194,13 +198,17 @@ struct ThreadWorkLoad {
inline size_t grid_dim(size_t i) const { return work_size[i]; }
};
/*! \brief Thread axis configuration */
-class ThreadAxisConfig {
+class LaunchParamConfig {
public:
- void Init(size_t base, const std::vector<std::string>& thread_axis_tags) {
+ void Init(size_t base, const std::vector<std::string>& launch_param_tags) {
base_ = base;
std::vector<bool> filled(6, false);
- for (size_t i = 0; i < thread_axis_tags.size(); ++i) {
- const std::string& tag = thread_axis_tags[i];
+ for (size_t i = 0; i < launch_param_tags.size(); ++i) {
+ const std::string& tag = launch_param_tags[i];
+ if (tag == kUseDynamicSharedMemoryTag) {
Review comment:
let us always assert this is the last tag that follows the arg_index map. Since it can be strange if the tag appears in the middle.
We can also handle a generalized case where the arg indicates the location of the parameter. In that case if the tag is
```[dyn_mem, threadIdx.x]```, then `x.values[base_+0]` would correspond to the dynamic shared mem size
--
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] masahi commented on a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r671690106
##########
File path: src/runtime/file_utils.cc
##########
@@ -64,12 +66,14 @@ void FunctionInfo::Save(dmlc::Stream* writer) const {
writer->Write(name);
writer->Write(arg_types);
writer->Write(thread_axis_tags);
+ writer->Write(use_dyn_shared_memory);
}
bool FunctionInfo::Load(dmlc::Stream* reader) {
if (!reader->Read(&name)) return false;
if (!reader->Read(&arg_types)) return false;
if (!reader->Read(&thread_axis_tags)) return false;
+ if (!reader->Read(&use_dyn_shared_memory)) return false;
Review comment:
@tqchen I didn't find a way to optionally read this attribute for the binary reader.
--
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 a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
tqchen commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r672230286
##########
File path: src/runtime/file_utils.cc
##########
@@ -64,12 +66,14 @@ void FunctionInfo::Save(dmlc::Stream* writer) const {
writer->Write(name);
writer->Write(arg_types);
writer->Write(thread_axis_tags);
+ writer->Write(use_dyn_shared_memory);
}
bool FunctionInfo::Load(dmlc::Stream* reader) {
if (!reader->Read(&name)) return false;
if (!reader->Read(&arg_types)) return false;
if (!reader->Read(&thread_axis_tags)) return false;
+ if (!reader->Read(&use_dyn_shared_memory)) return false;
Review comment:
I think. let us think about another way to do backward compact then.
Since binary backward compatibility is more important to us here.
Here is one solution that also helps future backward compatibility.
Let us rename `thread_axis_tags => launch_param_tags`, and use `dyn_shared_memory` as a tag for the launching parameter of dynamic shared memory. We can then change the `ThreadAxisConfig => LaunchParamConfig`.
##########
File path: src/runtime/file_utils.cc
##########
@@ -64,12 +66,14 @@ void FunctionInfo::Save(dmlc::Stream* writer) const {
writer->Write(name);
writer->Write(arg_types);
writer->Write(thread_axis_tags);
+ writer->Write(use_dyn_shared_memory);
}
bool FunctionInfo::Load(dmlc::Stream* reader) {
if (!reader->Read(&name)) return false;
if (!reader->Read(&arg_types)) return false;
if (!reader->Read(&thread_axis_tags)) return false;
+ if (!reader->Read(&use_dyn_shared_memory)) return false;
Review comment:
Let us try the ^ change and do a test to make sure we can load back old cuda binaries
--
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] masahi commented on a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r671689944
##########
File path: src/runtime/thread_storage_scope.h
##########
@@ -96,6 +98,8 @@ struct StorageScope {
return "global" + tag;
case StorageRank::kShared:
return "shared" + tag;
+ case StorageRank::kDynShared:
Review comment:
I changed dynamic shared memory to be a tagged memory for `kShared`, but I had to add a check `e->scope.tag != ".dyn"` in a few places in `lower_device_storage_access_info.cc` and `storage_rewrite.cc` to workaround error/segfault.
##########
File path: include/tvm/tir/function.h
##########
@@ -240,7 +240,7 @@ namespace attr {
*
* Call(f,
* [arg1, arg2, ..., arg_n,
- * work_size_1, work_size_2, ... work_size_m])
+ * work_size_1, work_size_2, ... work_size_m, dyn_shmem_size])
Review comment:
done
##########
File path: src/runtime/thread_storage_scope.h
##########
@@ -223,6 +232,7 @@ class ThreadAxisConfig {
w.work_size[arg_index_map_[i]] = size;
}
}
+ w.dyn_shmem_size = static_cast<size_t>(x.values[base_ + arg_index_map_.size()].v_int64);
Review comment:
done
--
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] yzh119 commented on a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by 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
[GitHub] [tvm] masahi commented on a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r673059490
##########
File path: src/runtime/file_utils.cc
##########
@@ -64,12 +66,14 @@ void FunctionInfo::Save(dmlc::Stream* writer) const {
writer->Write(name);
writer->Write(arg_types);
writer->Write(thread_axis_tags);
+ writer->Write(use_dyn_shared_memory);
}
bool FunctionInfo::Load(dmlc::Stream* reader) {
if (!reader->Read(&name)) return false;
if (!reader->Read(&arg_types)) return false;
if (!reader->Read(&thread_axis_tags)) return false;
+ if (!reader->Read(&use_dyn_shared_memory)) return false;
Review comment:
If we do the rename `thread_axis_tags => launch_param_tags`, the JSON reader would break backward compat at `helper.DeclareField("launch_param_tags", &launch_param_tags);` .
Or shall we do the rename but keep the old attribute name for JSON reader/writer? e.g.
`helper.DeclareField("thread_axis_tags", &launch_param_tags);`.
Or shall we have both new and old attribute names as optional field, e.g.
```
helper.DeclareOptionalField("launch_param_tags", &launch_param_tags);
helper.DeclareOptionalField("thread_axis_tags", &launch_param_tags); // for backward compatibility
```
--
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 #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
tqchen edited a comment on pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#issuecomment-881405667
Thanks @masahi I left some comments, also cc @Hzfengsy @vinx13 so they are also aware
--
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 a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
Hzfengsy commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r669724712
##########
File path: src/runtime/rocm/rocm_module.cc
##########
@@ -167,10 +167,12 @@ class ROCMWrappedFunc {
ThreadWorkLoad wl = thread_axis_cfg_.Extract(args);
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, packed_args, HIP_LAUNCH_PARAM_BUFFER_SIZE,
&packed_nbytes, HIP_LAUNCH_PARAM_END};
+ LOG(INFO) << "dynamic shared mem size: " << wl.dyn_shmem_size;
Review comment:
Please remove debug output
--
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] masahi commented on a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r673103103
##########
File path: src/runtime/file_utils.cc
##########
@@ -64,12 +66,14 @@ void FunctionInfo::Save(dmlc::Stream* writer) const {
writer->Write(name);
writer->Write(arg_types);
writer->Write(thread_axis_tags);
+ writer->Write(use_dyn_shared_memory);
}
bool FunctionInfo::Load(dmlc::Stream* reader) {
if (!reader->Read(&name)) return false;
if (!reader->Read(&arg_types)) return false;
if (!reader->Read(&thread_axis_tags)) return false;
+ if (!reader->Read(&use_dyn_shared_memory)) return false;
Review comment:
Went with the latter solution, and confirmed that an old saved module can be loaded via both json and binary.
--
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 a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
tqchen commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r671191780
##########
File path: src/runtime/thread_storage_scope.h
##########
@@ -96,6 +98,8 @@ struct StorageScope {
return "global" + tag;
case StorageRank::kShared:
return "shared" + tag;
+ case StorageRank::kDynShared:
Review comment:
Let us use the same storage scope as shared, but introduce a different tag. e.g. shared.dyn. The scope rank was used to rank the memory hierachy and needs to also made consistent with the thread rank. In this case dynamic shared memory is a special case of shared memory
--
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] masahi commented on a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r672217797
##########
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:
I realized that we must support this usage, since it is required, among others, for multiplying fp16 or int8 matrices and accumulate into fp32 or int32, for example (tensor core).
I'm thinking that we need to expose pointer cast in TIR, and somehow extend `Load` and `Store` node to be able to read/write with arbitrary pointers. I want to be able to express something like below :
```
extern __shared__ char buf[SIZE];
half* A_ptr = (half*)&buf[0];
half* B_ptr = (half*)&buf[128];
float32* C_ptr = (float32*)&buf[256];
...
C_ptr[0] = A_ptr[0] * B_ptr[0];
```
I didn't get what you meant by "translate to an intrinsic" @tqchen. Can we hide pointer cast and the load /stores that follow it behind intrinsics? For allocation and merging, I'm thinking about working exclusively with byte arrays and make use of `storage_rewrite` as is.
--
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 a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
tqchen commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r671191780
##########
File path: src/runtime/thread_storage_scope.h
##########
@@ -96,6 +98,8 @@ struct StorageScope {
return "global" + tag;
case StorageRank::kShared:
return "shared" + tag;
+ case StorageRank::kDynShared:
Review comment:
Let us use the same storage scope, but introduce a different tag. e.g. shared.dyn
##########
File path: src/runtime/thread_storage_scope.h
##########
@@ -96,6 +98,8 @@ struct StorageScope {
return "global" + tag;
case StorageRank::kShared:
return "shared" + tag;
+ case StorageRank::kDynShared:
Review comment:
Let us use the same storage scope as shared, but introduce a different tag. e.g. shared.dyn
--
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] yzh119 commented on a change in pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
yzh119 commented on a change in pull request #8466:
URL: https://github.com/apache/tvm/pull/8466#discussion_r670064809
##########
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:
@masahi this happens when you need to load both data and index into shared memory (e.g. sparse matrix multiplications).
--
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] masahi merged pull request #8466: [CUDA] Initial support for dynamic shared memory
Posted by GitBox <gi...@apache.org>.
masahi merged pull request #8466:
URL: https://github.com/apache/tvm/pull/8466
--
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