Port HIP translation to use ASR

In !839 (merged) it became too hard to keep the AST -> C++/HIP translation up to date after an AST refactoring so it was removed in !844 (merged); the translation needs to be ported to use ASR anyway and it wasn't actually used in production anyway. I created this issue to document the current state, what was working and how to re-enable it properly.

The last version in master that works is 737fc413. Here is how to use it, on the only two examples that were actually tested as part of the testsuite:

$ cat tests/subroutine4.f90 
subroutine triad(a, b, scalar, c)
real, intent(in) :: a(:), b(:), scalar
real, intent(out) :: c(:)
integer :: N, i
N = size(a)
do concurrent (i = 1:N)
    c(i) = a(i) + scalar * b(i)
end do
end subroutine

$ ./src/bin/cpptranslate --show-ast-cpp-hip tests/subroutine4.f90
#define blocksize 128

//Pass by value variables found in the loop body:  scalar
//Pass by reference variables found in the loop body:  b a c
__global__ void Tempkernelname(int N, float scalar, float *b, float *a, float *c){
    int i = blockIDx.x*blockDim.x+threadIdx.x;
    if (i >= N) return;
    c[i] = (a[i]) + ((scalar)*(b[i]));
}

void triad(float *a, size_t a_size, float *b, size_t b_size, float scalar, 
float *c, size_t c_size)
{
    size_t N;
    size_t i;
    N = a_size;
    int gridsize = (N + blocksize - 1)/blocksize;
    float *b_d;
    hipMalloc(&b_d, N*sizeof(float));
    hipMemcpy(b_d, b, N*sizeof(float), hipMemcpyHostToDevice);
    float *a_d;
    hipMalloc(&a_d, N*sizeof(float));
    hipMemcpy(a_d, a, N*sizeof(float), hipMemcpyHostToDevice);
    float *c_d;
    hipMalloc(&c_d, N*sizeof(float));
    hipMemcpy(c_d, c, N*sizeof(float), hipMemcpyHostToDevice);
    hipLaunchKernelGGL(Tempkernelname, dim3(gridsize), dim3(blocksize), 0, 0, N, scalar, b_d, a_d, c_d);

}

$ cat tests/subroutine6.f90 
subroutine triad(a, b, scalar, c)
real, intent(in) :: a(:), b(:), scalar
real, intent(out) :: c(:)
integer :: N, N2, i, j
N = size(a)
N2 = size(b)
do concurrent (i = 1:N)
    c(i) = a(i) + scalar * b(i)
end do

do concurrent (j = 1:N2)
    c(j) = b(j) + scalar
end do
end subroutine

$ ./src/bin/cpptranslate --show-ast-cpp-hip tests/subroutine6.f90
#define blocksize 128

//Pass by value variables found in the loop body:  scalar
//Pass by reference variables found in the loop body:  b c
__global__ void Tempkernelname1(int N2, float scalar, float *b, float *c){
    int j = blockIDx.x*blockDim.x+threadIdx.x;
    if (j >= N2) return;
    c[j] = (b[j]) + (scalar);
}

//Pass by value variables found in the loop body:  scalar
//Pass by reference variables found in the loop body:  b a c
__global__ void Tempkernelname(int N, float scalar, float *b, float *a, float *c){
    int i = blockIDx.x*blockDim.x+threadIdx.x;
    if (i >= N) return;
    c[i] = (a[i]) + ((scalar)*(b[i]));
}

void triad(float *a, size_t a_size, float *b, size_t b_size, float scalar, 
float *c, size_t c_size)
{
    size_t N;
    size_t N2;
    size_t i;
    size_t j;
    N = a_size;
    N2 = a_size;
    int gridsize = (N + blocksize - 1)/blocksize;
    float *b_d;
    hipMalloc(&b_d, N*sizeof(float));
    hipMemcpy(b_d, b, N*sizeof(float), hipMemcpyHostToDevice);
    float *a_d;
    hipMalloc(&a_d, N*sizeof(float));
    hipMemcpy(a_d, a, N*sizeof(float), hipMemcpyHostToDevice);
    float *c_d;
    hipMalloc(&c_d, N*sizeof(float));
    hipMemcpy(c_d, c, N*sizeof(float), hipMemcpyHostToDevice);
    hipLaunchKernelGGL(Tempkernelname, dim3(gridsize), dim3(blocksize), 0, 0, N, scalar, b_d, a_d, c_d);

    int gridsize1 = (N2 + blocksize - 1)/blocksize;
    float *b_d1;
    hipMalloc(&b_d1, N2*sizeof(float));
    hipMemcpy(b_d1, b, N2*sizeof(float), hipMemcpyHostToDevice);
    float *c_d1;
    hipMalloc(&c_d1, N2*sizeof(float));
    hipMemcpy(c_d1, c, N2*sizeof(float), hipMemcpyHostToDevice);
    hipLaunchKernelGGL(Tempkernelname1, dim3(gridsize1), dim3(blocksize), 0, 0, N2, scalar, b_d1, c_d1);

}

The translation was mainly implemented in the visit_DoConcurrentLoop() method here: https://gitlab.com/lfortran/lfortran/-/blob/737fc413dc4ac532cdf3db07f8ac9f06f1100e43/src/lfortran/ast_to_cpp_hip.cpp#L489.

What we should do to re-enable this feature is to take the new ASR based C++ backend, and enable to select how to parallelize arrays and loops. Currently Kokkos is hardwired in the C++ backend and we should also allow just raw C pointers and HIP, or any other combination that people would like to see.