Time loop for multiple GPU's

Question

I have a loop that has a similar structure as give below:

DO WHILE (t < tend)
  t = t + dt
  DO j = 1,dim

    var2 = 0
    CALL func_a(param1,var2)
    CALL MPI_COMM_RANK(comm(j),rank,ierr)
    CALL func_b(param1, var2,rank)

    var2 = var2 + 0.5d0*var3(j)
    CALL func_c(var3,param1)
    var2 = var2 - 0.5d0*var3

    CALL func_d(var2,params)
    DO ii = 1,dim
        CALL func_e(var2,param)
        CALL func_f(var2,param)
        CALL func_g(var1,var2,comm(j),t)
        var1 = var2   
    END DO
  END DO
END DO

This section is a time loop that I want to port to the GPU. All functions are parallel enough to be ported to the GPU individually.

All the variables named var1, var2, var3 are on the GPU memory copied to the GPU memory(devvar1, devvar2, devvar3) before the time loop starts. If I do the following:

DO WHILE (t < tend)
  t = t + dt
  DO j = 1,dim

    devvar2 = 0
    CALL cu_func_a<<<block,grid>>>(param1,devvar2)
    CALL cu_func_b<<<block,grid>>>(param1, devvar2)

    devvar2 = devvar2 + 0.5d0*devvar3(j)
    CALL cu_func_c<<<block,grid>>>(devvar3,param1)
    devvar2 = devvar2 - 0.5d0*devvar3

    CALL cu_func_d<<<block,grid>>>(devvar2,params)
    DO ii = 1,dim
        CALL cu_func_e<<<block,grid>>>(devvar2,param)
        CALL cu_func_f<<<block,grid>>>(devvar2,param)
        CALL cu_func_g<<<block,grid>>>(devvar1,devvar2,t)
        devvar1 = devvar2   
    END DO
  END DO
END DO

would it copy from the GPU memory to the GPU memory (in the same GPU) itself ? Or is there some unexpected behaviour possible.

Can I create another kernel (specifically to multiply, add and copy the vectors) and launch it within the kernels (cu_func_a, cu_func_b,...) so that this copying outside is not necessary.

Also I have another question. I am using multiple K20x GPU's, which have about 6 GB global memory. When I copy from the CPU to the GPU before the time loop, what happens if the overall memory I copy exceeds 6 GB ? Does it give me an error or I just get unexpected behaviour ?

Please let me know if anything above is unclear. I am happy to clarify and discuss ideas.

Thanks!

Answer 1

I divide the answer for the different question. (I answer to some)

Can I create another kernel (specifically to multiply, add and copy the vectors) and launch it within the kernels (cu_func_a, cu_func_b,...) so that this copying outside is not necessary.

Yes you can, this is a feature called dynamic parallelism. It was introduced in CUDA 5. In these links, part1 and part2 there is the explanation how to use with examples.

would it copy from the GPU memory to the GPU memory (in the same GPU) itself ?

I'm not sure I understand, if you question id if you can modify/copy a vector in the GPU memory yes you can. For example with Thrust you can copy (using '=') device vector in device vector or modify (see saxpy example) device vectors.

Note: there are different libraries for cuda for specific arguments, for optimum performance it is important to use them (see for example list, cups )

UPDATE

When I copy from the CPU to the GPU before the time loop, what happens if the overall memory I copy exceeds 6 GB ?

This depends by the GPU. From the CUDA 8 documentation, section J.1.3. GPU Memory Oversubscription, using the unified memory with:

• GPU < 6.x cite: "cannot allocate more managed memory than the physical size of GPU memory"
• GPU > 6.x cite: "in other words they can allocate, access, and share arrays larger than the total physical capacity of the system, enabling out-of-core processing of very large datasets. cudaMallocManaged will not run out of memory as long as there is enough system memory available for the allocation."