text: #include <hip/hip_runtime.h>
  #include "rpp_hip_host_decls.hpp"
  
  extern "C" __global__ void bitwise_AND(unsigned char *input1,
                                         unsigned char *input2,
                                         unsigned char *output,
                                         const unsigned int height,
                                         const unsigned int width,
                                         const unsigned int channel)
  {
      int id_x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
      int id_y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
      int id_z = hipBlockIdx_z * hipBlockDim_z + hipThreadIdx_z;
  
      if (id_x >= width || id_y >= height || id_z >= channel)
      {
          return;
      }
  
      int pixIdx = id_x + id_y * width + id_z * width * height;
      output[pixIdx] = input1[pixIdx] & input2[pixIdx];
  }
  
  extern "C" __global__ void bitwise_AND_batch(unsigned char *input1,
                                               unsigned char *input2,
                                               unsigned char *output,
                                               unsigned int *xroi_begin,
                                               unsigned int *xroi_end,
                                               unsigned int *yroi_begin,
                                               unsigned int *yroi_end,
                                               unsigned int *height,
                                               unsigned int *width,
                                               unsigned int *max_width,
                                               unsigned long long *batch_index,
                                               const unsigned int channel,
                                               unsigned int *inc, // use width * height for pln and 1 for pkd
                                               const int plnpkdindex) // use 1 pln 3 for pkd
  {
      int id_x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
      int id_y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
      int id_z = hipBlockIdx_z * hipBlockDim_z + hipThreadIdx_z;
  
      unsigned long pixIdx = 0;
  
      pixIdx = batch_index[id_z] + (id_x + id_y * max_width[id_z]) * plnpkdindex;
      if((id_y >= yroi_begin[id_z]) && (id_y <= yroi_end[id_z]) && (id_x >= xroi_begin[id_z]) && (id_x <= xroi_end[id_z]))
      {
          for(int indextmp = 0; indextmp < channel; indextmp++)
          {
              output[pixIdx] = input1[pixIdx] & input2[pixIdx];
              pixIdx += inc[id_z];
          }
      }
      else if((id_x < width[id_z]) && (id_y < height[id_z]))
      {
          for(int indextmp = 0; indextmp < channel; indextmp++)
          {
              output[pixIdx] = input1[pixIdx];
              pixIdx += inc[id_z];
          }
      }
  }
  
  RppStatus hip_exec_bitwise_AND_batch(Rpp8u *srcPtr1, Rpp8u *srcPtr2, Rpp8u *dstPtr, rpp::Handle& handle, RppiChnFormat chnFormat, Rpp32u channel, Rpp32s plnpkdind, Rpp32u max_height, Rpp32u max_width)
  {
      int localThreads_x = 32;
      int localThreads_y = 32;
      int localThreads_z = 1;
      int globalThreads_x = (max_width + 31) & ~31;
      int globalThreads_y = (max_height + 31) & ~31;
      int globalThreads_z = handle.GetBatchSize();
  
      hipLaunchKernelGGL(bitwise_AND_batch,
                         dim3(ceil((float)globalThreads_x/localThreads_x), ceil((float)globalThreads_y/localThreads_y), ceil((float)globalThreads_z/localThreads_z)),
                         dim3(localThreads_x, localThreads_y, localThreads_z),
                         0,
                         handle.GetStream(),
                         srcPtr1,
                         srcPtr2,
                         dstPtr,
                         handle.GetInitHandle()->mem.mgpu.roiPoints.x,
                         handle.GetInitHandle()->mem.mgpu.roiPoints.roiWidth,
                         handle.GetInitHandle()->mem.mgpu.roiPoints.y,
                         handle.GetInitHandle()->mem.mgpu.roiPoints.roiHeight,
                         handle.GetInitHandle()->mem.mgpu.srcSize.height,
                         handle.GetInitHandle()->mem.mgpu.srcSize.width,
                         handle.GetInitHandle()->mem.mgpu.maxSrcSize.width,
                         handle.GetInitHandle()->mem.mgpu.srcBatchIndex,
                         channel,
                         handle.GetInitHandle()->mem.mgpu.inc,
                         plnpkdind);
  
      return RPP_SUCCESS;
  }