CUDA kernel performance is bad using a for-loop : I have a 1D array in C-code that I would like to convert to a corresponding CUDA kernel that can execute with good performance on a GPU DEVICE. The architecture is Nvidia CUDA 10.x

I have tried the naive approach whereby the C-code is applied directly to a CUDA kernel and got correct results.

```CUDA

#include <cuda.h>
#include <cuComplex.h>
#include <stdio.h>

// DEVICE code:

__global__ void myKern(const int N, cuComplex *d_a, cuComplex *d_b){
  // global thread ID
  int i = blockIdx.x * blockDim.x + threadIdx.x;

  // check range
  if(i >= N) return;
 
  // *****************************************************
  // This appears to work, but is VERY bad performance
  // - can this for-loop be removed or defined differently
  //   for better performance ?
  for(; i < N; ++i){
    d_b[0].x *= d_a[i].x;
    d_b[0].y *= d_a[i].y;
  }
  // *****************************************************

}

// HOST code:

int main() {
  cuComplex *h_a, h_b;
  cuComplex *d_a, *d_b, chk;
  const int N = 5;

  h_b.x = 0.0f;
  h_b.y = 0.0f;

  // allocate HOST and DEVICE memory
  h_a = (cuComplex*)malloc(N*sizeof(cuComplex));
  cudaMalloc((void**)&d_a, N*sizeof(cuComplex));
  cudaMalloc((void**)&d_b, sizeof(cuComplex));

  // initialize input array
  for(int i = 0; i < N; ++i){
    h_a[i].x = (float)(i + 1);
    h_a[i].y = 1.0f
  }

  // copy to DEVICE
  cudaMemcpy(d_a, h_a, N*sizeof(cuComplex), cudaMemcpyHostToDevice);
  cudaMemcpy(d_b, &h_b, sizeof(cuComplex), cudaMemcpyHostToDevice);

  // call CUDA kernel
  int tpb = 1024;
  int nblocks = (N + tpb - 1) / tpb;
  myKern<<<nblocks,tpb>>>(N, d_a, d_b);

  // ******************************************
  // PROBLEM: I would like to convert this
  // for-loop to CUDA kernel that is optimized
  // for performance
  // ******************************************
  for(int i = 0; i < N; ++i){
    h_b.x *= h_a[i].x;
    h_b.y *= h_a[i].y;
  }
  // ******************************************

  // get results from DEVICE
  cudaMemcpy(&chk, d_b, sizeof(cuComplex), cudaMemcpyDeviceToHost);

  // compare HOST and DEVICE
  printf("HOST: %0.4f, DEVICE: %0.4f\n", h_b.x, chk.x);

  // free memory
  free(h_a);
  cudaFree(d_a); cudaFree(d_b);
```  

The results from HOST and DEVICE appear same, but performance is degraded on DEVICE due to the for-loop inside the CUDA kernel. I don't get an error and would just like to create a CUDA kernel that will properly execute the for-loop in C HOST code that has good performance for the DEVICE.