Skip to content

GitLab

Commit e536e096 authored by root's avatar root
Browse files

modified grouped gemm addressing method

parent aebd211c
Hide whitespace changes
Inline Side-by-side
Showing with 11 additions and 7 deletions
include/ck/tensor_operation/gpu/device/device_grouped_gemm_xdl.hpp
View file @ e536e096
...@@ -46,13 +46,17 @@ __global__ void ...@@ -46,13 +46,17 @@ __global__ void
const auto gemm_desc_ptr = const auto gemm_desc_ptr =
reinterpret_cast<const GemmDesc*>(cast_pointer_to_generic_address_space(gemm_descs_const)); reinterpret_cast<const GemmDesc*>(cast_pointer_to_generic_address_space(gemm_descs_const));
index_t group_id = 0; index_t left = 0;
for(index_t i = 0; i < group_count; i++) index_t right = group_count;
{ index_t group_id = index_t((left + right)/2);
group_id = while((!(block_id >= gemm_desc_ptr[group_id].BlockStart_ && block_id < gemm_desc_ptr[group_id].BlockEnd_)) && left <= right ){
(block_id >= gemm_desc_ptr[i].BlockStart_ && block_id < gemm_desc_ptr[i].BlockEnd_) if(block_id < gemm_desc_ptr[group_id].BlockStart_){
? i right = group_id;
: group_id; }
else{
left = group_id;
}
group_id = index_t((left + right)/2);
} }
GridwiseGemm::template Run<HasMainKBlockLoop>( GridwiseGemm::template Run<HasMainKBlockLoop>(
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment