! This file was generated by gpufort

 
! 
! Hints:
! Device variables in scope:
!       character                                   :: uplo, jobz

!       character                                   :: uplo, jobz

!       integer                                     :: n, nz, lda, lwork, istat, info

!       integer                                     :: n, nz, lda, lwork, istat, info

!       integer                                     :: n, nz, lda, lwork, istat, info

!       integer                                     :: n, nz, lda, lwork, istat, info

!       integer                                     :: n, nz, lda, lwork, istat, info

!       integer                                     :: n, nz, lda, lwork, istat, info

!       integer                                     :: lwork_h, liwork_h, ldz_h

!       integer                                     :: lwork_h, liwork_h, ldz_h

!       integer                                     :: lwork_h, liwork_h, ldz_h

!       integer                                     :: i, j, k, nb1, nb2, ib, mi, ldt, ldz, il, iu

!       integer                                     :: i, j, k, nb1, nb2, ib, mi, ldt, ldz, il, iu

!       integer                                     :: i, j, k, nb1, nb2, ib, mi, ldt, ldz, il, iu

!       integer                                     :: i, j, k, nb1, nb2, ib, mi, ldt, ldz, il, iu

!       integer                                     :: i, j, k, nb1, nb2, ib, mi, ldt, ldz, il, iu

!       integer                                     :: i, j, k, nb1, nb2, ib, mi, ldt, ldz, il, iu

!       integer                                     :: i, j, k, nb1, nb2, ib, mi, ldt, ldz, il, iu

!       integer                                     :: i, j, k, nb1, nb2, ib, mi, ldt, ldz, il, iu

!       integer                                     :: i, j, k, nb1, nb2, ib, mi, ldt, ldz, il, iu

!       integer                                     :: i, j, k, nb1, nb2, ib, mi, ldt, ldz, il, iu

!       integer                                     :: i, j, k, nb1, nb2, ib, mi, ldt, ldz, il, iu

!       integer                                     :: inde, indtau, indwrk, llwork, llwork_h, indwk2, indwk3, llwrk2

!       integer                                     :: inde, indtau, indwrk, llwork, llwork_h, indwk2, indwk3, llwrk2

!       integer                                     :: inde, indtau, indwrk, llwork, llwork_h, indwk2, indwk3, llwrk2

!       integer                                     :: inde, indtau, indwrk, llwork, llwork_h, indwk2, indwk3, llwrk2

!       integer                                     :: inde, indtau, indwrk, llwork, llwork_h, indwk2, indwk3, llwrk2

!       integer                                     :: inde, indtau, indwrk, llwork, llwork_h, indwk2, indwk3, llwrk2

!       integer                                     :: inde, indtau, indwrk, llwork, llwork_h, indwk2, indwk3, llwrk2

!       integer                                     :: inde, indtau, indwrk, llwork, llwork_h, indwk2, indwk3, llwrk2

!       real(8), parameter                          :: one = 1.0_8

!       type(dim3) :: blocks, threads

!       type(dim3) :: blocks, threads

function test_launch_krnl_e26a05_0_auto()
  ! errorCode > 0 implies that the test has failed
  use iso_c_binding
  use hip
  use dsyevd_gpu_kernels
  implicit none
  integer :: errorCode = 1
  ! TODO fix parameters
  ! - Add missing arguments
  ! - Determine size of arrays (typically indicated by 'type(c_ptr)' type)
  ! - Add target where we need a pointer
  integer(c_int),intent(IN) :: sharedMem
  type(c_ptr),value,intent(IN) :: stream
  integer,value :: n
  TODO declaration not found :: z
  TODO declaration not found :: a
  ! TODO Create initial data on host
  ! TODO Copy data to device ! (dest,src,size,direction)
  CALL hipCheck(hipMemcpy(???,c_loc(???),C_SIZEOF(???),hipMemcpyHostToDevice)) ! 
  CALL hipCheck(hipMemcpy(???,c_loc(???),C_SIZEOF(???),hipMemcpyHostToDevice)) ! 
  ! ... might be more (or less) than two memcopies 
  ! TODO run the test
  CALL launch_krnl_e26a05_0_auto(0,c_null_ptr,sharedMem,stream,n,z,a) ! Modify sharedMem if other than default 0
  CALL hipCheck(hipDeviceSynchronize())
  CALL launch_krnl_e26a05_0_cpu(0,c_null_ptr,sharedMem,stream,n,z,a)

  ! TODO Copy results back to host
  CALL hipCheck(hipMemcpy(c_loc(???),???,C_SIZEOF(???),hipMemcpyDeviceToHost)
  CALL hipCheck(hipMemcpy(c_loc(???),???,C_SIZEOF(???),hipMemcpyDeviceToHost)
  ! ... might be more (or less) than two memcopies 
  ! TODO Compare results
  ! TODO Update error code if the results do not match
  return errorCode
end function

 
! 
! Hints:
! Device variables in scope:
!       integer                               :: n, k, ldv, ldt, ldw

!       integer                               :: n, k, ldv, ldt, ldw

!       integer                               :: n, k, ldv, ldt, ldw

!       integer                               :: n, k, ldv, ldt, ldw

!       integer                               :: n, k, ldv, ldt, ldw

!       real(8), dimension(ldv, k), device    :: v

!       real(8), dimension(k), device         :: tau

!       real(8), dimension(ldt, k), device    :: t

!       real(8), dimension(ldw, k), device    :: w

!       integer                               :: i, j, istat

!       integer                               :: i, j, istat

!       integer                               :: i, j, istat

!       type(dim3)                            :: threads

function test_launch_krnl_b1f342_1_auto()
  ! errorCode > 0 implies that the test has failed
  use iso_c_binding
  use hip
  use dsyevd_gpu_kernels
  implicit none
  integer :: errorCode = 1
  ! TODO fix parameters
  ! - Add missing arguments
  ! - Determine size of arrays (typically indicated by 'type(c_ptr)' type)
  ! - Add target where we need a pointer
  integer(c_int),intent(IN) :: sharedMem
  type(c_ptr),value,intent(IN) :: stream
  type(c_ptr),value :: w
  integer,value :: n
  integer,value :: k
  type(c_ptr),value :: v
  ! TODO Create initial data on host
  ! TODO Copy data to device ! (dest,src,size,direction)
  CALL hipCheck(hipMemcpy(???,c_loc(???),C_SIZEOF(???),hipMemcpyHostToDevice)) ! 
  CALL hipCheck(hipMemcpy(???,c_loc(???),C_SIZEOF(???),hipMemcpyHostToDevice)) ! 
  ! ... might be more (or less) than two memcopies 
  ! TODO run the test
  CALL launch_krnl_b1f342_1_auto(0,c_null_ptr,sharedMem,stream,w,n,k,v) ! Modify sharedMem if other than default 0
  CALL hipCheck(hipDeviceSynchronize())
  CALL launch_krnl_b1f342_1_cpu(0,c_null_ptr,sharedMem,stream,w,n,k,v)

  ! TODO Copy results back to host
  CALL hipCheck(hipMemcpy(c_loc(???),???,C_SIZEOF(???),hipMemcpyDeviceToHost)
  CALL hipCheck(hipMemcpy(c_loc(???),???,C_SIZEOF(???),hipMemcpyDeviceToHost)
  ! ... might be more (or less) than two memcopies 
  ! TODO Compare results
  ! TODO Update error code if the results do not match
  return errorCode
end function

 
! 
! Hints:
! Device variables in scope:
!       integer                               :: m, n, k, ldv, ldt, ldc, ldw, ldwork, istat

!       integer                               :: m, n, k, ldv, ldt, ldc, ldw, ldwork, istat

!       integer                               :: m, n, k, ldv, ldt, ldc, ldw, ldwork, istat

!       integer                               :: m, n, k, ldv, ldt, ldc, ldw, ldwork, istat

!       integer                               :: m, n, k, ldv, ldt, ldc, ldw, ldwork, istat

!       integer                               :: m, n, k, ldv, ldt, ldc, ldw, ldwork, istat

!       integer                               :: m, n, k, ldv, ldt, ldc, ldw, ldwork, istat

!       integer                               :: m, n, k, ldv, ldt, ldc, ldw, ldwork, istat

!       integer                               :: m, n, k, ldv, ldt, ldc, ldw, ldwork, istat

!       integer                               :: i, j

!       integer                               :: i, j

!       real(8), dimension(ldv, k), device    :: v

!       real(8), dimension(ldt, k), device    :: t

!       real(8), dimension(ldw, k), device    :: w

!       real(8), dimension(ldc, n), device    :: c

!       real(8), dimension(ldwork, k), device :: work

function test_launch_krnl_b95769_2_auto()
  ! errorCode > 0 implies that the test has failed
  use iso_c_binding
  use hip
  use dsyevd_gpu_kernels
  implicit none
  integer :: errorCode = 1
  ! TODO fix parameters
  ! - Add missing arguments
  ! - Determine size of arrays (typically indicated by 'type(c_ptr)' type)
  ! - Add target where we need a pointer
  integer(c_int),intent(IN) :: sharedMem
  type(c_ptr),value,intent(IN) :: stream
  integer,value :: m
  type(c_ptr),value :: w
  integer,value :: k
  type(c_ptr),value :: v
  ! TODO Create initial data on host
  ! TODO Copy data to device ! (dest,src,size,direction)
  CALL hipCheck(hipMemcpy(???,c_loc(???),C_SIZEOF(???),hipMemcpyHostToDevice)) ! 
  CALL hipCheck(hipMemcpy(???,c_loc(???),C_SIZEOF(???),hipMemcpyHostToDevice)) ! 
  ! ... might be more (or less) than two memcopies 
  ! TODO run the test
  CALL launch_krnl_b95769_2_auto(0,c_null_ptr,sharedMem,stream,m,w,k,v) ! Modify sharedMem if other than default 0
  CALL hipCheck(hipDeviceSynchronize())
  CALL launch_krnl_b95769_2_cpu(0,c_null_ptr,sharedMem,stream,m,w,k,v)

  ! TODO Copy results back to host
  CALL hipCheck(hipMemcpy(c_loc(???),???,C_SIZEOF(???),hipMemcpyDeviceToHost)
  CALL hipCheck(hipMemcpy(c_loc(???),???,C_SIZEOF(???),hipMemcpyDeviceToHost)
  ! ... might be more (or less) than two memcopies 
  ! TODO Compare results
  ! TODO Update error code if the results do not match
  return errorCode
end function

 
! Fortran implementation:
!     attributes(global) subroutine finish_T_block_kernel(N, T, ldt, tau)
!        implicit none
!        integer, value                     :: N, ldt
!        real(8), dimension(ldt, K), device :: T
!        real(8), dimension(K), device      :: tau
!        ! T_s contains only lower triangular elements of T in linear array, by row
!        real(8), dimension(2080), shared   :: T_s
!        ! (i,j) --> ((i-1)*i/2 + j)
! #define IJ2TRI(i,j) (ISHFT((i-1)*i,-1) + j)
! 
!        integer     :: tid, tx, ty, i, j, k, diag
!        complex(8)  :: cv
! 
!        tx = threadIdx%x
!        ty = threadIdx%y
!        tid = (threadIdx%y - 1)*blockDim%x + tx ! Linear thread id
! 
!        ! Load T into shared memory
!        if (tx <= N) then
!           do j = ty, N, blockDim%y
!              cv = tau(j)
!              if (tx > j) then
!                 T_s(IJ2TRI(tx, j)) = -cv*T(tx, j)
!              else if (tx == j) then
!                 T_s(IJ2TRI(tx, j)) = cv
!              endif
!           end do
!        end if
! 
!        call syncthreads()
! 
!        ! Perform column by column update by first thread column
!        do i = N - 1, 1, -1
!           if (ty == 1) then
!              if (tx > i .and. tx <= N) then
!                 cv = 0.0d0
!                 do j = i + 1, tx
!                    cv = cv + T_s(IJ2TRI(j, i))*T_s(IJ2TRI(tx, j))
!                 end do
!              endif
! 
!           endif
! 
!           call syncthreads()
!           if (ty == 1 .and. tx > i .and. tx <= N) then
!              T_s(IJ2TRI(tx, i)) = cv
!           endif
!           call syncthreads()
! 
!        end do
! 
!        call syncthreads()
! 
!        ! Write T_s to global
!        if (tx <= N) then
!           do j = ty, N, blockDim%y
!              if (tx >= j) then
!                 T(tx, j) = T_s(IJ2TRI(tx, j))
!              endif
!           end do
!        end if
! 
!     end subroutine finish_T_block_kernel
! 
function test_launch_finish_t_block_kernel()
  ! errorCode > 0 implies that the test has failed
  use iso_c_binding
  use hip
  use dsyevd_gpu_kernels
  implicit none
  integer :: errorCode = 1
  ! TODO fix parameters
  ! - Add missing arguments
  ! - Determine size of arrays (typically indicated by 'type(c_ptr)' type)
  ! - Add target where we need a pointer
  type(dim3,,intent(IN, :: grid
  type(dim3,,intent(IN, :: block
  integer(c_int,,intent(IN, :: sharedMem
  type(c_ptr,,value,intent(IN, :: stream
  integer,value :: n
  integer,value :: ldt
  type(c_ptr),value :: _t
  type(c_ptr),value :: _tau
  ! TODO Create initial data on host
  ! TODO Copy data to device ! (dest,src,size,direction)
  CALL hipCheck(hipMemcpy(???,c_loc(???),C_SIZEOF(???),hipMemcpyHostToDevice)) ! 
  CALL hipCheck(hipMemcpy(???,c_loc(???),C_SIZEOF(???),hipMemcpyHostToDevice)) ! 
  ! ... might be more (or less) than two memcopies 
  ! TODO run the test
  CALL launch_finish_t_block_kernel(0,c_null_ptr,grid,block,sharedMem,stream,n,ldt,t,tau) ! Modify sharedMem if other than default 0
  CALL hipCheck(hipDeviceSynchronize())
  CALL launch_finish_t_block_kernel_cpu(0,c_null_ptr,grid,block,sharedMem,stream,n,ldt,t,tau)

  ! TODO Copy results back to host
  CALL hipCheck(hipMemcpy(c_loc(???),???,C_SIZEOF(???),hipMemcpyDeviceToHost)
  CALL hipCheck(hipMemcpy(c_loc(???),???,C_SIZEOF(???),hipMemcpyDeviceToHost)
  ! ... might be more (or less) than two memcopies 
  ! TODO Compare results
  ! TODO Update error code if the results do not match
  return errorCode
end function


program test_dsyevd_gpu_kernels
  implicit none
  integer :: globalErrorCode = 0, errorCode, fails = 0, tests = 0
  ! declare test functions and return type
  integer :: test_launch_krnl_e26a05_0_auto
  integer :: test_launch_krnl_b1f342_1_auto
  integer :: test_launch_krnl_b95769_2_auto
  integer :: test_launch_finish_t_block_kernel
  write(*,*) "SUITE test_dsyevd_gpu_kernels run ..."
  errorCode = test_launch_krnl_e26a05_0_auto()
  IF (errorCode > 0) THEN
    fails = fails + 1
    write(*,*) "TEST test_launch_krnl_e26a05_0_auto ... FAILURE"
  ELSE 
    write(*,*) "TEST test_launch_krnl_e26a05_0_auto ... SUCCESS"
  END IF
  tests = tests + 1
  globalErrorCode = globalErrorCode + errorCode
  errorCode = test_launch_krnl_b1f342_1_auto()
  IF (errorCode > 0) THEN
    fails = fails + 1
    write(*,*) "TEST test_launch_krnl_b1f342_1_auto ... FAILURE"
  ELSE 
    write(*,*) "TEST test_launch_krnl_b1f342_1_auto ... SUCCESS"
  END IF
  tests = tests + 1
  globalErrorCode = globalErrorCode + errorCode
  errorCode = test_launch_krnl_b95769_2_auto()
  IF (errorCode > 0) THEN
    fails = fails + 1
    write(*,*) "TEST test_launch_krnl_b95769_2_auto ... FAILURE"
  ELSE 
    write(*,*) "TEST test_launch_krnl_b95769_2_auto ... SUCCESS"
  END IF
  tests = tests + 1
  globalErrorCode = globalErrorCode + errorCode
  errorCode = test_launch_finish_t_block_kernel()
  IF (errorCode > 0) THEN
    fails = fails + 1
    write(*,*) "TEST test_launch_finish_t_block_kernel ... FAILURE"
  ELSE 
    write(*,*) "TEST test_launch_finish_t_block_kernel ... SUCCESS"
  END IF
  tests = tests + 1
  globalErrorCode = globalErrorCode + errorCode

  IF (globalErrorCode > 0) THEN
    write(*,*) "SUITE test_dsyevd_gpu_kernels ... FAILURE passed:",(tests-fails)," failed:",fails," total:",tests
  ELSE 
    write(*,*) "SUITE test_dsyevd_gpu_kernels ... SUCCESS passed:",(tests-fails)," failed:",fails," total:",tests
  END IF
end program test_dsyevd_gpu_kernels