Skip to content

GitLab

Commit e83c7061 authored by root's avatar root
Browse files

modified addressing method in device_grouped_gemm_xdl.hpp

parent e536e096
Hide whitespace changes
Inline Side-by-side
Showing with 12 additions and 7 deletions
include/ck/tensor_operation/gpu/device/device_grouped_gemm_xdl.hpp
View file @ e83c7061
...@@ -46,17 +46,22 @@ __global__ void ...@@ -46,17 +46,22 @@ __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 left = 0; index_t left = 0;
index_t right = group_count; index_t right = group_count;
index_t group_id = index_t((left + right)/2); index_t group_id = index_t((left + right) / 2);
while((!(block_id >= gemm_desc_ptr[group_id].BlockStart_ && block_id < gemm_desc_ptr[group_id].BlockEnd_)) && left <= right ){ while((!(block_id >= gemm_desc_ptr[group_id].BlockStart_ &&
if(block_id < gemm_desc_ptr[group_id].BlockStart_){ block_id < gemm_desc_ptr[group_id].BlockEnd_)) &&
left <= right)
{
if(block_id < gemm_desc_ptr[group_id].BlockStart_)
{
right = group_id; right = group_id;
} }
else{ else
{
left = group_id; left = group_id;
} }
group_id = index_t((left + right)/2); 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