hip refactor for reduce, reduction, replicate, reshape and reverse

[Bob-Chen222]

Description of changes:
refactor for reduce, reduction, replicate, reshape and reverse

Related Issues:

Linked Issues:

• Issue Hip kernel refactor #1296

Issues closed by this PR:

• Closes #

---

This change is 

[codecov]

Codecov Report

All modified and coverable lines are covered by tests ✅

Project coverage is 38.10%. Comparing base (6fe5dba) to head (d785ece).

Additional details and impacted files

@@              Coverage Diff               @@
##           repo-refactor    #1403   +/-   ##
==============================================
  Coverage          38.10%   38.10%           
==============================================
  Files                167      167           
  Lines               5026     5026           
  Branches             246      246           
==============================================
  Hits                1915     1915           
  Misses              3111     3111           

Flag	Coverage Δ
unittests	38.10% <ø> (ø)

Flags with carried forward coverage won't be shown. Click here to find out more.

[reyna-abhyankar]

Reviewable status: 0 of 5 files reviewed, 8 unresolved discussions (waiting on @Bob-Chen222)

---

lib/kernels/src/hip/reduce_kernels.cpp line 27 at r1 (raw file):

                                 OperatorType const &op_type,
                                 size_t const &reduction_size,
                                 ArrayShape input_shape,

ArrayShape const & for both

---

lib/kernels/src/hip/reduce_kernels.cpp line 92 at r1 (raw file):

      // When the output is the average of multiple input elements
      // we need to scale the gradients by 1.0 / reduction_size
      alpha = 1.0f / m->reduction_size;

m.reduction_size

---

lib/kernels/src/hip/reduce_kernels.cpp line 97 at r1 (raw file):

      assert(false);
  }
  checkCUDNN(miopenOpTensor(m.handle.dnn,

should be AddTensor, see the cuda file

---

lib/kernels/src/hip/replicate_kernels.cpp line 45 at r1 (raw file):

    checkCUDA(hipMemcpyAsync(input.get<T>(),
                             output.get<T>(),
                             input.shape.num_elements() * sizeof(T),

... * size_of_datatype(T)

---

lib/kernels/src/hip/replicate_kernels.cpp line 58 at r1 (raw file):

                  size_t num_replicas) {
    size_t total_elements = input.shape.num_elements() * num_replicas;
    hipLaunchKernelGGL(HIP_KERNEL_NAME(replicate_backward_kernel<T>),

replicate_backward_kernel<real_type<T>>
since T is of DataType which is just a wrapper around the actual data type.

---

lib/kernels/src/hip/reshape_kernels.cpp line 37 at r1 (raw file):

    checkCUDA(hipMemcpyAsync(output.get<T>(),
                             input.get<T>(),
                             input.shape.num_elements() * sizeof(T),

See size_of_datatype comment above

---

lib/kernels/src/hip/reshape_kernels.cpp line 50 at r1 (raw file):

                  GenericTensorAccessorR const &output) {
    float alpha = 1.0f;
    hipLaunchKernelGGL(HIP_KERNEL_NAME(apply_add_with_scale<T>),

See real_type comment above

---

lib/kernels/src/hip/reshape_kernels.cpp line 58 at r1 (raw file):

                       output.get<T>(),
                       input.shape.num_elements(),
                       (T)alpha);

static_cast<real_type<T>>(alpha). The only time you can just use T here is for input.get<T> because GenericTensorAccessorW has a template function that takes in a DataType and spits out the pointer cast to the underlying data type. I think I fixed all of these in the cuda version so just double check that stuff.

[Bob-Chen222]

Reviewable status: 0 of 5 files reviewed, 8 unresolved discussions (waiting on @reyna-abhyankar)

---

lib/kernels/src/hip/reduce_kernels.cpp line 27 at r1 (raw file):

Previously, reyna-abhyankar (Reyna Abhyankar) wrote…

ArrayShape const & for both

Done.

---

lib/kernels/src/hip/reduce_kernels.cpp line 92 at r1 (raw file):

Previously, reyna-abhyankar (Reyna Abhyankar) wrote…

m.reduction_size

Done.

---

lib/kernels/src/hip/reduce_kernels.cpp line 97 at r1 (raw file):

Previously, reyna-abhyankar (Reyna Abhyankar) wrote…

should be AddTensor, see the cuda file

Done.

---

lib/kernels/src/hip/replicate_kernels.cpp line 45 at r1 (raw file):

Previously, reyna-abhyankar (Reyna Abhyankar) wrote…

... * size_of_datatype(T)

Done.

---

lib/kernels/src/hip/replicate_kernels.cpp line 58 at r1 (raw file):

Previously, reyna-abhyankar (Reyna Abhyankar) wrote…

replicate_backward_kernel<real_type<T>>
since T is of DataType which is just a wrapper around the actual data type.

Done.

---

lib/kernels/src/hip/reshape_kernels.cpp line 37 at r1 (raw file):

Previously, reyna-abhyankar (Reyna Abhyankar) wrote…

See size_of_datatype comment above

Done.

---

lib/kernels/src/hip/reshape_kernels.cpp line 50 at r1 (raw file):

Previously, reyna-abhyankar (Reyna Abhyankar) wrote…

See real_type comment above

Done.

---

lib/kernels/src/hip/reshape_kernels.cpp line 58 at r1 (raw file):

Previously, reyna-abhyankar (Reyna Abhyankar) wrote…

static_cast<real_type<T>>(alpha). The only time you can just use T here is for input.get<T> because GenericTensorAccessorW has a template function that takes in a DataType and spits out the pointer cast to the underlying data type. I think I fixed all of these in the cuda version so just double check that stuff.

Done.

[reyna-abhyankar]

Reviewable status: 0 of 5 files reviewed, 1 unresolved discussion (waiting on @Bob-Chen222)

---

lib/kernels/src/hip/reduce_kernels.cpp line 97 at r1 (raw file):

Previously, Bob-Chen222 (Bob Chen) wrote…

Done.

Actually, this is my mistake. It should be miopenOpTensor, I just didn't see that you had already passed in the miopenTensorOpAdd for the op type. You can revert back to the previous function call (but still keep m. instead of m->)

[Bob-Chen222]

Reviewable status: 0 of 5 files reviewed, 1 unresolved discussion (waiting on @reyna-abhyankar)

---

lib/kernels/src/hip/reduce_kernels.cpp line 97 at r1 (raw file):

Previously, reyna-abhyankar (Reyna Abhyankar) wrote…

Actually, this is my mistake. It should be miopenOpTensor, I just didn't see that you had already passed in the miopenTensorOpAdd for the op type. You can revert back to the previous function call (but still keep m. instead of m->)

Done.