Commit 633199a1 authored by Shucai Xiao's avatar Shucai Xiao
Browse files

fix cppcheck error.

parent 65702b8a
...@@ -42,14 +42,14 @@ void logsoftmax(hipStream_t stream, const argument& result, const argument& arg, ...@@ -42,14 +42,14 @@ void logsoftmax(hipStream_t stream, const argument& result, const argument& arg,
size_t blk_idx = idx.group; size_t blk_idx = idx.group;
using type = device_type<std::remove_cv_t<typename decltype(output)::value_type>>; using type = device_type<std::remove_cv_t<typename decltype(output)::value_type>>;
MIGRAPHX_DEVICE_SHARED type lds_data[max_block_size + 2]; MIGRAPHX_DEVICE_SHARED type lds_data[max_block_size + 1];
auto batch_idx = desc_batch.multi(blk_idx); auto batch_idx = desc_batch.multi(blk_idx);
auto data_idx = batch_idx; auto data_idx = batch_idx;
// load data to lds and compute the batch max // load data to lds and compute the batch max
size_t remaining_item_num = batch_item_num; size_t remaining_item_num = batch_item_num;
size_t thread_num = (batch_item_num + block_size - 1) / block_size * block_size; size_t round_item_num = (batch_item_num + block_size - 1) / block_size * block_size;
lds_data[block_size] = input_ptr[0]; lds_data[block_size] = input_ptr[0];
for(size_t i = thr_idx; i < thread_num; i += block_size) for(size_t i = thr_idx; i < round_item_num; i += block_size)
{ {
if(i < batch_item_num) if(i < batch_item_num)
{ {
...@@ -70,7 +70,7 @@ void logsoftmax(hipStream_t stream, const argument& result, const argument& arg, ...@@ -70,7 +70,7 @@ void logsoftmax(hipStream_t stream, const argument& result, const argument& arg,
lds_data[block_size] = 0; lds_data[block_size] = 0;
remaining_item_num = batch_item_num; remaining_item_num = batch_item_num;
for(size_t i = thr_idx; i < thread_num; i += block_size) for(size_t i = thr_idx; i < round_item_num; i += block_size)
{ {
if(i < batch_item_num) if(i < batch_item_num)
{ {
......
...@@ -11,7 +11,7 @@ namespace gpu { ...@@ -11,7 +11,7 @@ namespace gpu {
namespace device { namespace device {
template <class T> template <class T>
__device__ void reduce_max(T* data_ptr, size_t block_size, size_t thr_idx, size_t item_num) inline __device__ void reduce_max(T* data_ptr, size_t block_size, size_t thr_idx, size_t item_num)
{ {
auto stride = (item_num + 1) / 2; auto stride = (item_num + 1) / 2;
while(true) while(true)
...@@ -39,7 +39,7 @@ __device__ void reduce_max(T* data_ptr, size_t block_size, size_t thr_idx, size_ ...@@ -39,7 +39,7 @@ __device__ void reduce_max(T* data_ptr, size_t block_size, size_t thr_idx, size_
} }
template <class T> template <class T>
__device__ void reduce_sum(T* data_ptr, size_t block_size, size_t thr_idx, size_t item_num) inline __device__ void reduce_sum(T* data_ptr, size_t block_size, size_t thr_idx, size_t item_num)
{ {
auto stride = (item_num + 1) / 2; auto stride = (item_num + 1) / 2;
while(true) while(true)
......
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