Commit a0934d4e authored by Andreas Marek's avatar Andreas Marek
Browse files

Merge branch 'master' into ELPA_GPU

parents 98a4db33 33a94bfc
......@@ -17,7 +17,12 @@ libelpa@SUFFIX@_la_SOURCES = src/mod_precision.f90 \
src/check_for_gpu.F90 \
src/mod_cuda.F90 \
src/interface_c_kernel.F90 \
src/elpa2_compute.F90 \
src/mod_pack_unpack_real.F90 \
src/elpa2_kernels/mod_single_hh_trafo_real.F90 \
src/mod_compute_hh_trafo_real.F90 \
src/mod_compute_hh_trafo_complex.F90 \
src/mod_pack_unpack_complex.F90 \
src/elpa2_compute.F90 \
src/elpa2.F90 \
src/elpa_c_interface.F90 \
src/elpa_qr/qr_utils.f90 \
......@@ -314,6 +319,9 @@ elpa2.i: $(top_srcdir)/src/elpa2.F90
elpa1.i: $(top_srcdir)/src/elpa1.F90
$(CPP) $(CPPFLAGS) -I$(top_builddir)/ -I$(top_srcdir)/ -c $(top_srcdir)/src/elpa1.F90 -o $@
mod_compute_hh_trafo_real.i: $(top_srcdir)/src/mod_compute_hh_trafo_real.F90
$(CPP) $(CPPFLAGS) -I$(top_builddir)/ -c $(top_srcdir)/src/mod_compute_hh_trafo_real.F90 -o $@
include doxygen.am
CLEANFILES = \
......
......@@ -298,6 +298,8 @@ function solve_evp_real_1stage(na, nev, a, lda, ev, q, ldq, nblk, matrixCols, mp
integer(kind=ik), intent(in) :: na, nev, lda, ldq, nblk, matrixCols, mpi_comm_rows, mpi_comm_cols
real(kind=rk) :: a(lda,matrixCols), ev(na), q(ldq,matrixCols)
! was
! real a(lda,*), q(ldq,*)
integer(kind=ik) :: my_prow, my_pcol, mpierr
real(kind=rk), allocatable :: e(:), tau(:)
......@@ -397,6 +399,8 @@ function solve_evp_complex_1stage(na, nev, a, lda, ev, q, ldq, nblk, matrixCols,
integer(kind=ik), intent(in) :: na, nev, lda, ldq, nblk, matrixCols, mpi_comm_rows, mpi_comm_cols
complex(kind=ck) :: a(lda,matrixCols), q(ldq,matrixCols)
! was
! complex a(lda,*), q(ldq,*)
real(kind=rk) :: ev(na)
integer(kind=ik) :: my_prow, my_pcol, np_rows, np_cols, mpierr
......
......@@ -138,6 +138,8 @@ module ELPA1_compute
integer(kind=ik) :: na, lda, nblk, matrixCols, mpi_comm_rows, mpi_comm_cols
real(kind=rk) :: a(lda,matrixCols), d(na), e(na), tau(na)
! was
! real a(lda,*)
integer(kind=ik), parameter :: max_stored_rows = 32
......@@ -479,6 +481,8 @@ module ELPA1_compute
integer(kind=ik) :: na, nqc, lda, ldq, nblk, matrixCols, mpi_comm_rows, mpi_comm_cols
real(kind=rk) :: a(lda,matrixCols), q(ldq,matrixCols), tau(na)
! was
! real a(lda,*), q(ldq,*)
integer(kind=ik) :: max_stored_rows
......@@ -911,6 +915,8 @@ module ELPA1_compute
integer(kind=ik) :: na, lda, nblk, matrixCols, mpi_comm_rows, mpi_comm_cols
complex(kind=ck) :: a(lda,matrixCols), tau(na)
! was
! complex a(lda,*)
real(kind=rk) :: d(na), e(na)
integer(kind=ik), parameter :: max_stored_rows = 32
......@@ -1278,6 +1284,8 @@ module ELPA1_compute
integer(kind=ik) :: na, nqc, lda, ldq, nblk, matrixCols, mpi_comm_rows, mpi_comm_cols
complex(kind=ck) :: a(lda,matrixCols), q(ldq,matrixCols), tau(na)
! was
! complex a(lda,*), q(ldq,*)
integer(kind=ik) :: max_stored_rows
......@@ -1678,6 +1686,8 @@ module ELPA1_compute
integer(kind=ik) :: na, nev, ldq, nblk, matrixCols, mpi_comm_rows, mpi_comm_cols
real(kind=rk) :: d(na), e(na), q(ldq,matrixCols)
! was
! real q(ldq,*)
integer(kind=ik) :: i, j, n, np, nc, nev1, l_cols, l_rows
integer(kind=ik) :: my_prow, my_pcol, np_rows, np_cols, mpierr
......@@ -1902,6 +1912,8 @@ module ELPA1_compute
integer(kind=ik) :: na, nev, nqoff, ldq, nblk, matrixCols, mpi_comm_rows
real(kind=rk) :: d(na), e(na), q(ldq,matrixCols)
! was
! real q(ldq,*)
integer(kind=ik), parameter :: min_submatrix_size = 16 ! Minimum size of the submatrices to be used
......@@ -2175,6 +2187,8 @@ module ELPA1_compute
mpi_comm_cols, npc_0, npc_n
integer(kind=ik) :: l_col(na), p_col(na), l_col_out(na), p_col_out(na)
real(kind=rk) :: d(na), e, q(ldq,matrixCols)
! was
! real q(ldq,*)
integer(kind=ik), parameter :: max_strip=128
......@@ -3309,6 +3323,8 @@ module ELPA1_compute
integer(kind=ik) :: na, lda, nblk, matrixCols, mpi_comm_rows, mpi_comm_cols
real(kind=rk) :: a(lda,matrixCols)
! was
! real a(lda, *)
integer(kind=ik) :: my_prow, my_pcol, np_rows, np_cols, mpierr
integer(kind=ik) :: l_cols, l_rows, l_col1, l_row1, l_colx, l_rowx
......@@ -3490,6 +3506,8 @@ module ELPA1_compute
integer(kind=ik) :: na, lda, nblk, matrixCols, mpi_comm_rows, mpi_comm_cols
real(kind=rk) :: a(lda,matrixCols)
! was
! real a(lda,*)
integer(kind=ik) :: my_prow, my_pcol, np_rows, np_cols, mpierr
integer(kind=ik) :: l_cols, l_rows, l_col1, l_row1, l_colx, l_rowx
......@@ -3626,6 +3644,8 @@ module ELPA1_compute
integer(kind=ik) :: na, lda, nblk, matrixCols, mpi_comm_rows, mpi_comm_cols
complex(kind=ck) :: a(lda,matrixCols)
!was
! complex a(lda,*)
integer(kind=ik) :: my_prow, my_pcol, np_rows, np_cols, mpierr
integer(kind=ik) :: l_cols, l_rows, l_col1, l_row1, l_colx, l_rowx
......@@ -3803,6 +3823,8 @@ module ELPA1_compute
integer(kind=ik) :: na, lda, nblk, matrixCols, mpi_comm_rows, mpi_comm_cols
complex(kind=ck) :: a(lda,matrixCols)
! was
! complex a(lda,*)
integer(kind=ik) :: my_prow, my_pcol, np_rows, np_cols, mpierr
integer(kind=ik) :: l_cols, l_rows, l_col1, l_row1, l_colx, l_rowx
......
......@@ -156,6 +156,8 @@ contains
mpi_comm_cols, mpi_comm_all
integer(kind=ik), intent(in) :: nblk
real(kind=rk), intent(inout) :: a(lda,matrixCols), ev(na), q(ldq,matrixCols)
! was
! real a(lda,*), q(ldq,*)
real(kind=rk), allocatable :: hh_trans_real(:,:)
integer(kind=ik) :: my_pe, n_pes, my_prow, my_pcol, np_rows, np_cols, mpierr
......@@ -429,6 +431,8 @@ function solve_evp_complex_2stage(na, nev, a, lda, ev, q, ldq, nblk, &
integer(kind=ik) :: THIS_COMPLEX_ELPA_KERNEL
integer(kind=ik), intent(in) :: na, nev, lda, ldq, nblk, matrixCols, mpi_comm_rows, mpi_comm_cols, mpi_comm_all
complex(kind=ck), intent(inout) :: a(lda,matrixCols), q(ldq,matrixCols)
! was
! complex a(lda,*), q(ldq,*)
real(kind=rk), intent(inout) :: ev(na)
complex(kind=ck), allocatable :: hh_trans_complex(:,:)
......
......@@ -143,6 +143,8 @@ module ELPA2_compute
integer(kind=ik) :: na, lda, nblk, nbw, matrixCols, numBlocks, mpi_comm_rows, mpi_comm_cols
real(kind=rk) :: a(lda,matrixCols), tmat(nbw,nbw,numBlocks)
! was
! real a(lda,*), tmat(nbw,nbw,*)
real(kind=rk) :: eps
logical, intent(in) :: useGPU
integer(kind=ik) :: my_prow, my_pcol, np_rows, np_cols, mpierr
......@@ -625,9 +627,7 @@ module ELPA2_compute
else
if (l_rows>0) &
call dsyrk('U','T',n_cols,l_rows,1.d0,vmrCPU,ubound(vmrCPU,dim=1),0.d0,vav,ubound(vav,dim=1))
endif
call symm_matrix_allreduce(n_cols,vav, nbw, nbw,mpi_comm_rows)
! Calculate triangular matrix T for block Householder Transformation
......@@ -1238,6 +1238,8 @@ module ELPA2_compute
integer(kind=ik) :: na, nqc, lda, ldq, nblk, nbw, matrixCols, numBlocks, mpi_comm_rows, mpi_comm_cols
real(kind=rk) :: a(lda,matrixCols), q(ldq,matrixCols), tmat(nbw, nbw, numBlocks)
! was
! real a(lda,*), q(ldq,*), tmat(nbw,nbw,*)
integer(kind=ik) :: my_prow, my_pcol, np_rows, np_cols, mpierr
integer(kind=ik) :: max_blocks_row, max_blocks_col, max_local_rows, &
......@@ -1689,6 +1691,8 @@ module ELPA2_compute
integer(kind=ik), intent(in) :: na, nb, nblk, lda, matrixCols, mpi_comm_rows, mpi_comm_cols, mpi_comm
real(kind=rk), intent(in) :: a(lda,matrixCols)
! was
! real a(lda,*)
real(kind=rk), intent(out) :: d(na), e(na) ! set only on PE 0
real(kind=rk), intent(out), &
allocatable :: hh_trans_real(:,:)
......@@ -2470,12 +2474,16 @@ module ELPA2_compute
#endif
use cuda_functions
use precision
use pack_unpack_real
use compute_hh_trafo_real
implicit none
logical, intent(in) :: useGPU
integer(kind=ik), intent(in) :: THIS_REAL_ELPA_KERNEL
integer(kind=ik), intent(in) :: na, nev, nblk, nbw, ldq, matrixCols, mpi_comm_rows, mpi_comm_cols
real(kind=rk) :: q(ldq,matrixCols)
! was
! real q(ldq,*)
real(kind=rk), intent(out) :: hh_trans_real(:,:)
integer(kind=ik) :: np_rows, my_prow, np_cols, my_pcol
......@@ -2776,7 +2784,8 @@ module ELPA2_compute
!$omp parallel do private(my_thread), schedule(static, 1)
do my_thread = 1, max_threads
call unpack_row_real_cpu_openmp(row,i-limits(ip),my_thread)
call unpack_row_real_cpu_openmp(a, row,i-limits(ip),my_thread, stripe_count, &
thread_width, stripe_width, l_nev)
enddo
!$omp end parallel do
#ifdef HAVE_DETAILED_TIMINGS
......@@ -2791,7 +2800,7 @@ module ELPA2_compute
call MPI_Recv(row_group(:, row_group_size), l_nev, MPI_REAL8, src, 0, mpi_comm_rows, MPI_STATUS_IGNORE, mpierr)
else
call MPI_Recv(row, l_nev, MPI_REAL8, src, 0, mpi_comm_rows, MPI_STATUS_IGNORE, mpierr)
call unpack_row_real_cpu(row,i-limits(ip))
call unpack_row_real_cpu(a, row,i-limits(ip), stripe_count, stripe_width, last_stripe_width)
endif
#endif /* WITH_OPENMP */
......@@ -2812,7 +2821,8 @@ module ELPA2_compute
!$omp parallel do private(my_thread), schedule(static, 1)
do my_thread = 1, max_threads
call unpack_row_real_cpu_openmp(row,i-limits(ip),my_thread)
call unpack_row_real_cpu_openmp(a, row,i-limits(ip),my_thread, &
stripe_count, thread_width, stripe_width, l_nev)
enddo
!$omp end parallel do
#ifdef HAVE_DETAILED_TIMINGS
......@@ -2820,13 +2830,12 @@ module ELPA2_compute
#endif
#else /* WITH_OPENMP */
if (useGPU) then
! An unpacking of the current row group may occur before queuing the next row
call unpack_and_prepare_row_group_real_gpu(i - limits(ip), .false.)
row_group(:, row_group_size) = q(src_offset, 1:l_nev)
else
call unpack_row_real_cpu(row,i-limits(ip))
call unpack_row_real_cpu(a, row,i-limits(ip), stripe_count, stripe_width, last_stripe_width)
endif
#endif /* WITH_OPENMP */
endif
......@@ -2869,7 +2878,8 @@ module ELPA2_compute
!$omp parallel do private(my_thread), schedule(static, 1)
do my_thread = 1, max_threads
call unpack_row_real_cpu_openmp(row,i-limits(my_prow),my_thread)
call unpack_row_real_cpu_openmp(a, row,i-limits(my_prow),my_thread, &
stripe_count, thread_width, stripe_width, l_nev)
enddo
!$omp end parallel do
#ifdef HAVE_DETAILED_TIMINGS
......@@ -2877,14 +2887,13 @@ module ELPA2_compute
#endif
#else /* WITH_OPENMP */
if (useGPU) then
! An unpacking of the current row group may occur before queuing the next row
call unpack_and_prepare_row_group_real_gpu(i - limits(my_prow), .false.)
call MPI_Recv(row_group(:, row_group_size), l_nev, MPI_REAL8, src, 0, mpi_comm_rows, MPI_STATUS_IGNORE, mpierr)
else
call MPI_Recv(row, l_nev, MPI_REAL8, src, 0, mpi_comm_rows, MPI_STATUS_IGNORE, mpierr)
call unpack_row_real_cpu(row,i-limits(my_prow))
call unpack_row_real_cpu(a, row,i-limits(my_prow), stripe_count, stripe_width, last_stripe_width)
endif
#endif /* WITH_OPENMP */
......@@ -3084,7 +3093,6 @@ module ELPA2_compute
current_tv_off = 0 ! Offset of next row to be broadcast
! ------------------- start of work loop -------------------
a_off = 0 ! offset in A (to avoid unnecessary shifts)
......@@ -3300,8 +3308,10 @@ module ELPA2_compute
a(1:csw,a_off+1:a_off+top_msg_length,i,my_thread) = &
reshape(top_border_recv_buffer(b_off+1:b_off+b_len,i), (/ csw, top_msg_length /))
endif
call compute_hh_trafo_real(0, current_local_n, i, my_thread, &
THIS_REAL_ELPA_KERNEL)
call compute_hh_trafo_real_cpu_openmp(a,stripe_width,a_dim2,stripe_count, max_threads, &
a_off, nbw, max_blk_size, bcast_buffer, kernel_flops, kernel_time, &
0, current_local_n, i, my_thread, &
THIS_REAL_ELPA_KERNEL)
enddo
!$omp end parallel do
#ifdef HAVE_DETAILED_TIMINGS
......@@ -3309,8 +3319,10 @@ module ELPA2_compute
#endif
#else /* WITH_OPENMP */
call compute_hh_trafo_real(0, current_local_n, i, &
THIS_REAL_ELPA_KERNEL)
call compute_hh_trafo_real_cpu(a, stripe_width,a_dim2,stripe_count, &
a_off, nbw, max_blk_size, bcast_buffer, kernel_flops, kernel_time, &
0, current_local_n, i, &
last_stripe_width, THIS_REAL_ELPA_KERNEL)
#endif /* WITH_OPENMP */
!send_b
......@@ -3365,8 +3377,10 @@ module ELPA2_compute
!$omp parallel do private(my_thread, b_len, b_off), schedule(static, 1)
do my_thread = 1, max_threads
call compute_hh_trafo_real(current_local_n - bottom_msg_length, bottom_msg_length, i, my_thread, &
THIS_REAL_ELPA_KERNEL)
call compute_hh_trafo_real_cpu_openmp(a, stripe_width,a_dim2,stripe_count, max_threads, &
a_off, nbw, max_blk_size, bcast_buffer, kernel_flops, kernel_time, &
current_local_n - bottom_msg_length, bottom_msg_length, i, my_thread, &
THIS_REAL_ELPA_KERNEL)
enddo
!$omp end parallel do
#ifdef HAVE_DETAILED_TIMINGS
......@@ -3384,8 +3398,10 @@ module ELPA2_compute
top_recv_tag, mpi_comm_rows, bottom_send_request(i), mpierr)
endif
#else /* WITH_OPENMP */
call compute_hh_trafo_real(current_local_n - bottom_msg_length, bottom_msg_length, i, &
THIS_REAL_ELPA_KERNEL)
call compute_hh_trafo_real_cpu(a, stripe_width,a_dim2,stripe_count, &
a_off, nbw, max_blk_size, bcast_buffer, kernel_flops, kernel_time, &
current_local_n - bottom_msg_length, bottom_msg_length, i, &
last_stripe_width, THIS_REAL_ELPA_KERNEL)
!send_b
call MPI_Wait(bottom_send_request(i), MPI_STATUS_IGNORE, mpierr)
......@@ -3421,8 +3437,10 @@ module ELPA2_compute
!$omp parallel do private(my_thread), schedule(static, 1)
do my_thread = 1, max_threads
call compute_hh_trafo_real(top_msg_length, current_local_n-top_msg_length-bottom_msg_length, i, my_thread, &
THIS_REAL_ELPA_KERNEL)
call compute_hh_trafo_real_cpu_openmp(a,stripe_width,a_dim2,stripe_count, max_threads, &
a_off, nbw, max_blk_size, bcast_buffer, kernel_flops, kernel_time, &
top_msg_length, current_local_n-top_msg_length-bottom_msg_length, i, my_thread, &
THIS_REAL_ELPA_KERNEL)
enddo
!$omp end parallel do
#ifdef HAVE_DETAILED_TIMINGS
......@@ -3430,9 +3448,10 @@ module ELPA2_compute
#endif
#else /* WITH_OPENMP */
call compute_hh_trafo_real(top_msg_length, current_local_n-top_msg_length-bottom_msg_length, i, &
THIS_REAL_ELPA_KERNEL)
call compute_hh_trafo_real_cpu(a, stripe_width,a_dim2,stripe_count, &
a_off, nbw, max_blk_size, bcast_buffer, kernel_flops, kernel_time, &
top_msg_length, current_local_n-top_msg_length-bottom_msg_length, i, &
last_stripe_width, THIS_REAL_ELPA_KERNEL)
#endif /* WITH_OPENMP */
!wait_t
......@@ -3480,7 +3499,9 @@ module ELPA2_compute
a(1:csw,a_off+1:a_off+top_msg_length,i,my_thread) = &
reshape(top_border_recv_buffer(b_off+1:b_off+b_len,i), (/ csw, top_msg_length /))
endif
call compute_hh_trafo_real(0, top_msg_length, i, my_thread, THIS_REAL_ELPA_KERNEL)
call compute_hh_trafo_real_cpu_openmp(a, stripe_width,a_dim2,stripe_count, max_threads, &
a_off, nbw, max_blk_size, bcast_buffer, kernel_flops, kernel_time, &
0, top_msg_length, i, my_thread, THIS_REAL_ELPA_KERNEL)
enddo
!$omp end parallel do
#ifdef HAVE_DETAILED_TIMINGS
......@@ -3488,7 +3509,10 @@ module ELPA2_compute
#endif
#else /* WITH_OPENMP */
call compute_hh_trafo_real(0, top_msg_length, i, THIS_REAL_ELPA_KERNEL)
call compute_hh_trafo_real_cpu(a, stripe_width,a_dim2,stripe_count, &
a_off, nbw, max_blk_size, bcast_buffer, kernel_flops, kernel_time, &
0, top_msg_length, i, &
last_stripe_width, THIS_REAL_ELPA_KERNEL)
#endif /* WITH_OPENMP */
endif
......@@ -3637,7 +3661,7 @@ module ELPA2_compute
call pack_row_group_real_gpu(result_buffer(:, :, nbuf), j * nblk + a_off, nblk)
else
do i = 1, nblk
call pack_row_real_cpu(result_buffer(:,i,nbuf),j*nblk+i+a_off)
call pack_row_real_cpu(a, result_buffer(:,i,nbuf),j*nblk+i+a_off, stripe_width, last_stripe_width, stripe_count)
enddo
endif
call MPI_Isend(result_buffer(1,1,nbuf), l_nev*nblk, MPI_REAL8, dst, &
......@@ -3941,102 +3965,6 @@ module ELPA2_compute
return
contains
subroutine pack_row_real_cpu(row, n)
#ifdef HAVE_DETAILED_TIMINGS
use timings
#endif
use precision
implicit none
real(kind=rk) :: row(:)
integer(kind=ik) :: n, i, noff, nl
#ifdef WITH_OPENMP
integer(kind=ik) :: nt
#endif
#ifdef HAVE_DETAILED_TIMINGS
call timer%start("pack_row_real_cpu")
#endif
#ifdef WITH_OPENMP
do nt = 1, max_threads
do i = 1, stripe_count
noff = (nt-1)*thread_width + (i-1)*stripe_width
nl = min(stripe_width, nt*thread_width-noff, l_nev-noff)
if (nl<=0) exit
row(noff+1:noff+nl) = a(1:nl,n,i,nt)
enddo
enddo
#else
do i=1,stripe_count
nl = merge(stripe_width, last_stripe_width, i<stripe_count)
noff = (i-1)*stripe_width
row(noff+1:noff+nl) = a(1:nl,n,i)
enddo
#endif
#ifdef HAVE_DETAILED_TIMINGS
call timer%stop("pack_row_real_cpu")
#endif
end subroutine pack_row_real_cpu
#ifdef WITH_OPENMP
subroutine unpack_row_real_cpu_openmp(row, n, my_thread)
#ifdef HAVE_DETAILED_TIMINGS
use timings
#endif
use precision
implicit none
! Private variables in OMP regions (my_thread) should better be in the argument list!
integer(kind=ik), intent(in) :: n, my_thread
real(kind=rk), intent(in) :: row(:)
integer(kind=ik) :: i, noff, nl
#ifdef HAVE_DETAILED_TIMINGS
call timer%start("unpack_row_real_cpu_openmp")
#endif
do i=1,stripe_count
noff = (my_thread-1)*thread_width + (i-1)*stripe_width
nl = min(stripe_width, my_thread*thread_width-noff, l_nev-noff)
if(nl<=0) exit
a(1:nl,n,i,my_thread) = row(noff+1:noff+nl)
enddo
#ifdef HAVE_DETAILED_TIMINGS
call timer%stop("unpack_row_real_cpu_openmp")
#endif
end subroutine unpack_row_real_cpu_openmp
#else /* WITH_OPENMP */
subroutine unpack_row_real_cpu(row, n)
#ifdef HAVE_DETAILED_TIMINGS
use timings
#endif
use precision
implicit none
real(kind=rk) :: row(:)
integer(kind=ik) :: n, i, noff, nl
#ifdef HAVE_DETAILED_TIMINGS
call timer%start("unpack_row_real_cpu")
#endif
do i=1,stripe_count
nl = merge(stripe_width, last_stripe_width, i<stripe_count)
noff = (i-1)*stripe_width
a(1:nl,n,i) = row(noff+1:noff+nl)
enddo
#ifdef HAVE_DETAILED_TIMINGS
call timer%stop("unpack_row_real_cpu")
#endif
end subroutine unpack_row_real_cpu
#endif /* WITH_OPENMP */
! Pack a filled row group (i.e. an array of consecutive rows)
subroutine pack_row_group_real_gpu(rows, n_offset, row_count)
use cuda_c_kernel
......@@ -4197,397 +4125,8 @@ module ELPA2_compute
! - "compute_hh_trafo_kernel" is the Fortran equivalent of the C kernel
! - "compute_hh_trafo_single_kernel" is the reference Fortran kernel
#ifdef WITH_OPENMP
subroutine compute_hh_trafo_real_openmp(off, ncols, istripe, my_thread, THIS_REAL_ELPA_KERNEL)
#else
subroutine compute_hh_trafo_real(off, ncols, istripe, THIS_REAL_ELPA_KERNEL)
#endif
use precision
#if defined(WITH_REAL_GENERIC_SIMPLE_KERNEL)
use real_generic_simple_kernel, only : double_hh_trafo_generic_simple
#endif
!#if defined(WITH_REAL_GENERIC_KERNEL)
! use real_generic_kernel, only : double_hh_trafo_generic
!#endif
#if defined(WITH_REAL_BGP_KERNEL)
use real_bgp_kernel, only : double_hh_trafo_bgp
#endif
#if defined(WITH_REAL_BGQ_KERNEL)
use real_bgq_kernel, only : double_hh_trafo_bgq
#endif
#ifdef HAVE_DETAILED_TIMINGS
use timings
#endif
use cuda_c_kernel
implicit none
integer(kind=ik), intent(in) :: THIS_REAL_ELPA_KERNEL
! Private variables in OMP regions (my_thread) should better be in the argument list!
integer(kind=ik) :: off, ncols, istripe
#ifdef WITH_OPENMP
integer(kind=ik) :: my_thread, noff
#endif
integer(kind=ik) :: j, nl, jj, jjj
real(kind=rk) :: w(nbw,6), ttt
if (THIS_REAL_ELPA_KERNEL .eq. REAL_ELPA_KERNEL_GPU) then
! ncols - indicates the number of HH reflectors to apply; at least 1 must be available
if (ncols < 1) return
endif
#ifdef HAVE_DETAILED_TIMINGS
#ifdef WITH_OPENMP
call timer%start("compute_hh_trafo_real_openmp")
#else
call timer%start("compute_hh_trafo_real")
#endif
#endif
ttt = mpi_wtime()
#ifndef WITH_OPENMP
nl = merge(stripe_width, last_stripe_width, istripe<stripe_count)
#else
if (THIS_REAL_ELPA_KERNEL .eq. REAL_ELPA_KERNEL_GPU) then
print *,"compute_hh_trafo_real: not yet implemented"
stop
endif
if (istripe<stripe_count) then
nl = stripe_width
else
noff = (my_thread-1)*thread_width + (istripe-1)*stripe_width
nl = min(my_thread*thread_width-noff, l_nev-noff)
if (nl<=0) return
endif
#endif
if (THIS_REAL_ELPA_KERNEL .eq. REAL_ELPA_KERNEL_GPU) then
dev_offset = (0 + (a_off * stripe_width) + ( (istripe - 1) * stripe_width *a_dim2 )) *size_of_real_datatype
call launch_compute_hh_trafo_c_kernel_real(a_dev + dev_offset, bcast_buffer_dev, hh_dot_dev, &
hh_tau_dev, nl, nbw, stripe_width, off, ncols)
else ! not CUDA kernel
#if defined(WITH_NO_SPECIFIC_REAL_KERNEL)
if (THIS_REAL_ELPA_KERNEL .eq. REAL_ELPA_KERNEL_AVX_BLOCK2 .or. &
THIS_REAL_ELPA_KERNEL .eq. REAL_ELPA_KERNEL_GENERIC .or. &
THIS_REAL_ELPA_KERNEL .eq. REAL_ELPA_KERNEL_GENERIC_SIMPLE .or. &
THIS_REAL_ELPA_KERNEL .eq. REAL_ELPA_KERNEL_SSE .or. &
THIS_REAL_ELPA_KERNEL .eq. REAL_ELPA_KERNEL_BGP .or. &
THIS_REAL_ELPA_KERNEL .eq. REAL_ELPA_KERNEL_BGQ) then
#endif /* WITH_NO_SPECIFIC_REAL_KERNEL */
!FORTRAN CODE / X86 INRINISIC CODE / BG ASSEMBLER USING 2 HOUSEHOLDER VECTORS
#if defined(WITH_REAL_GENERIC_KERNEL)
#if defined(WITH_NO_SPECIFIC_REAL_KERNEL)
if (THIS_REAL_ELPA_KERNEL .eq. REAL_ELPA_KERNEL_GENERIC) then
#endif /* WITH_NO_SPECIFIC_REAL_KERNEL */
do j = ncols, 2, -2
w(:,1) = bcast_buffer(1:nbw,j+off)
w(:,2) = bcast_buffer(1:nbw,j+off-1)
#ifdef WITH_OPENMP
call double_hh_trafo_generic(a(1,j+off+a_off-1,istripe,my_thread), w, &