Commit a3aebe12 authored by Giovanni La Mura's avatar Giovanni La Mura
Browse files

Introduce optional fully offloaded matrix calculation

parent bef40b91

src/cluster/cluster.cpp

+23 −25
Original line number Diff line number Diff line
@@ -894,15 +894,12 @@ int cluster_jxi488_cycle( 
    {

      magma_cms(vec_am, cid->c1, cid->proc_device); // Initialize uninverted matrix on GPU

      

#pragma omp target data use_device_ptr(vec_am) device(cid->proc_device)

      {

      magma_zinvert_resident(vec_am, (magma_int_t)ndit, jer, queue, cid->proc_device, rs);

      magma_queue_sync(queue);

      }

      if (jer == 0) {

	// Get the final inverted matrix.

#pragma omp target update from(vec_am[0:ndit*ndit]) device(cid->proc_device)

      } else {

      

      magma_ztm(vec_am, cid->c1, cid->proc_device);

      

      if (jer != 0) {

	sprintf(virtual_line, "ERROR: matrix inversion returned code %d!\n", jer);

	message = virtual_line;

	logger->err(message);
@@ -913,6 +910,7 @@ int cluster_jxi488_cycle( 
    // TODO: implement full offload pipeline without MAGMA

    cms_flat(cid->am[0], cid->c1);

    invert_matrix(cid->am, ndit, jer, output_path, jxi488, mxndm, cid->proc_device, rs);

    ztm(cid->am, cid->c1);

#endif // USE_MAGMA

  } else {

#ifdef USE_NVTX
@@ -971,7 +969,6 @@ int cluster_jxi488_cycle( 
      return jer;

      // break; // jxi488 loop: goes to memory clean

    }

  }

    interval_start = chrono::high_resolution_clock::now();

#ifdef USE_NVTX

    nvtxRangePush("Average calculation");
@@ -990,6 +987,7 @@ int cluster_jxi488_cycle( 
    outam3->write_to_disk(outam3_name);

    delete outam3;

#endif // DEBUG_AM

  }

  if (idfc >= 0) {

    if (jxi488 == jwtm) {

      int nlemt = 2 * cid->c1->nlem;