Port HIP translation to use ASR

Original issue: https://gitlab.com/lfortran/lfortran/-/issues/323

In https://gitlab.com/lfortran/lfortran/-/merge_requests/839 it became too hard to keep the AST -> C++/HIP translation up to date after an AST refactoring so it was removed in https://gitlab.com/lfortran/lfortran/-/merge_requests/844; 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 737fc413dc4ac532cdf3db07f8ac9f06f1100e43. 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.

Start by reading the ASR-based C++ backend and the historical implementation in src/lfortran/ast_to_cpp_hip.cpp, especially visit_DoConcurrentLoop(). Run the examples in tests/subroutine4.f90 and tests/subroutine6.f90 using the historical command as a reference. Done means the ASR backend can select raw C pointers and HIP, alongside Kokkos, and the tested translations work again.