Merge branch 'release-2.1.0-rc14'

.cp2k/Makefile

@@ -178,7 +178,7 @@ OBJ_SRC_FILES += $(shell cd $(SRCDIR); find . -name "*.c")
 # if compiling with GPU acceleration
 ifneq ($(ACC),)
   # All *.cpp files belong to the accelerator backend
-  OBJ_SRC_FILES += $(shell cd $(SRCDIR); find . ! -name "acc_cuda.cpp" ! -name "acc_hip.cpp" -name "*.cpp")
+  OBJ_SRC_FILES += $(shell cd $(SRCDIR); find . ! -name "acc_cuda.cpp" ! -name "acc_hip.cpp" ! -name "hipblas.cpp" -name "*.cpp")
   # if compiling with nvcc
   ifneq (,$(findstring nvcc,$(ACC)))
     OBJ_SRC_FILES += $(LIBSMM_ACC_DIR)/../cuda/acc_cuda.cpp
@@ -187,6 +187,7 @@ ifneq ($(ACC),)
   # if compiling with hipcc
   else ifneq (,$(findstring hipcc,$(ACC)))
     OBJ_SRC_FILES += $(LIBSMM_ACC_DIR)/../hip/acc_hip.cpp
+    OBJ_SRC_FILES += $(LIBSMM_ACC_DIR)/../hipblaswrap/hipblas.cpp
   endif
 endif
 

VERSION

@@ -1,6 +1,6 @@
 MAJOR = 2
 MINOR = 1
-PATCH = 0-rc13
+PATCH = 0-rc14
 # A specific DATE (YYYY-MM-DD) fixes an official release, otherwise
 # it is considered Development version.
 DATE  =

src/acc/acc_stream.cpp

@@ -54,7 +54,7 @@ extern "C" int acc_stream_create(void** stream_p, const char* name, int priority
   if (acc_error_check(ACC(GetLastError)())) return -1;
 
 #ifdef __CUDA_PROFILING
-  nvtxNameCudaStreamA(*custream, name);
+  nvtxNameCudaStreamA(*acc_stream, name);
 #endif
 
     return 0;

src/acc/cuda/dbcsr_cuda_nvtx_cu.cu

@@ -31,7 +31,7 @@ const uint32_t colormap[] = { 0xFFFFFF00,  // Yellow
 //==============================================================================
 extern "C" int cuda_nvtx_range_push_cu(const char* message) {
 
-    //assembling event attribute
+    // assembling event attribute
     nvtxEventAttributes_t eventAttrib = {0};
     eventAttrib.version = NVTX_VERSION;
     eventAttrib.size = NVTX_EVENT_ATTRIB_STRUCT_SIZE;
@@ -45,7 +45,7 @@ extern "C" int cuda_nvtx_range_push_cu(const char* message) {
     eventAttrib.colorType = NVTX_COLOR_ARGB;
     eventAttrib.color = colormap[hash%14];
 
-    //these field could be fild with useful stuff
+    // these field could be filled with useful stuff
     eventAttrib.payloadType = NVTX_PAYLOAD_TYPE_INT64;
     eventAttrib.payload.llValue = 123;
     eventAttrib.category = 42;

src/dbcsr_api.F

@@ -936,8 +936,12 @@ SUBROUTINE dbcsr_create_template(matrix, name, template, &
 
       INTEGER, DIMENSION(:), POINTER, CONTIGUOUS         :: cont_row_blk_size, cont_col_blk_size
 
+      IF (PRESENT(row_blk_size) .NEQV. PRESENT(col_blk_size)) THEN
+         DBCSR_ABORT("Both row_blk_size and col_blk_size must be provided!")
+      ENDIF
+
       ! Make the array contiguous, avoid to change API
-      IF (PRESENT(row_blk_size) .AND. PRESENT(col_blk_size)) THEN
+      IF (PRESENT(row_blk_size)) THEN
          ! Avoid to change API
          ALLOCATE (cont_row_blk_size(SIZE(row_blk_size)), cont_col_blk_size(SIZE(col_blk_size)))
          cont_row_blk_size(:) = row_blk_size(:)
@@ -951,7 +955,7 @@ SUBROUTINE dbcsr_create_template(matrix, name, template, &
       ENDIF
 
       IF (PRESENT(dist)) THEN
-         IF (PRESENT(row_blk_size) .AND. PRESENT(col_blk_size)) THEN
+         IF (PRESENT(row_blk_size)) THEN
             CALL dbcsr_create_prv(matrix%prv, template%prv, name, &
                                   dist%prv, matrix_type, &
                                   row_blk_size=cont_row_blk_size, col_blk_size=cont_col_blk_size, &
@@ -966,7 +970,7 @@ SUBROUTINE dbcsr_create_template(matrix, name, template, &
                                   replication_type=replication_type)
          ENDIF
       ELSE
-         IF (PRESENT(row_blk_size) .AND. PRESENT(col_blk_size)) THEN
+         IF (PRESENT(row_blk_size)) THEN
             CALL dbcsr_create_prv(matrix%prv, template%prv, name, &
                                   matrix_type=matrix_type, &
                                   row_blk_size=cont_row_blk_size, col_blk_size=cont_col_blk_size, &

src/tensors/dbcsr_tensor.F

@@ -93,7 +93,6 @@ MODULE dbcsr_tensor
       dbcsr_t_reserve_blocks, &
       dbcsr_t_copy_matrix_to_tensor, &
       dbcsr_t_copy_tensor_to_matrix, &
-      dbcsr_t_need_contract, &
       dbcsr_t_contract_index, &
       dbcsr_t_batched_contract_init, &
       dbcsr_t_batched_contract_finalize
@@ -598,6 +597,7 @@ SUBROUTINE dbcsr_t_contract_expert(alpha, tensor_1, tensor_2, beta, tensor_3, &
 
       IF (PRESENT(flop)) flop = 0
       IF (PRESENT(result_index)) result_index = 0
+      IF (PRESENT(nblks_local)) nblks_local = 0
 
       IF (PRESENT(move_data)) THEN
          move_data_1 = move_data
@@ -1626,61 +1626,6 @@ SUBROUTINE dbcsr_t_contract_index(alpha, tensor_1, tensor_2, beta, tensor_3, &
                                    result_index=result_index)
    END SUBROUTINE
 
-   FUNCTION dbcsr_t_need_contract(tensor_1, tensor_2, contract_1, notcontract_1, &
-                                  contract_2, notcontract_2, &
-                                  bounds_1, bounds_2, bounds_3)
-      !! check whether we need to contract two tensors (whether the contraction indices of both tensors
-      !! have common occupied blocks). See documentation of dbcsr_t_contract.
-
-      TYPE(dbcsr_t_type), INTENT(INOUT)  :: tensor_1, tensor_2
-      INTEGER, DIMENSION(:), INTENT(IN)  :: contract_1, contract_2, &
-                                            notcontract_1, notcontract_2
-      INTEGER, DIMENSION(2, SIZE(contract_1)), &
-         OPTIONAL                                    :: bounds_1
-      INTEGER, DIMENSION(2, SIZE(notcontract_1)), &
-         OPTIONAL                                    :: bounds_2
-      INTEGER, DIMENSION(2, SIZE(notcontract_2)), &
-         OPTIONAL                                    :: bounds_3
-      INTEGER :: i
-      LOGICAL, DIMENSION(:), ALLOCATABLE :: ind1, ind2
-      INTEGER, DIMENSION(2, ndims_tensor(tensor_1)) :: bounds_t1
-      INTEGER, DIMENSION(2, ndims_tensor(tensor_2)) :: bounds_t2
-      LOGICAL :: dbcsr_t_need_contract
-      CHARACTER(LEN=*), PARAMETER :: routineN = 'dbcsr_t_need_contract', &
-                                     routineP = moduleN//':'//routineN
-      INTEGER :: handle, mp_comm
-
-      CALL timeset(routineN, handle)
-
-      mp_comm = tensor_1%pgrid%mp_comm_2d
-
-      CALL dbcsr_t_map_bounds_to_tensors(tensor_1, tensor_2, &
-                                         contract_1, notcontract_1, &
-                                         contract_2, notcontract_2, &
-                                         bounds_t1, bounds_t2, &
-                                         bounds_1=bounds_1, bounds_2=bounds_2, bounds_3=bounds_3)
-
-      dbcsr_t_need_contract = .TRUE.
-
-      DO i = 1, SIZE(contract_1)
-         CALL dbcsr_t_get_nonzero_index(tensor_1, contract_1(i), bounds_t1, ind1)
-         CALL dbcsr_t_get_nonzero_index(tensor_2, contract_2(i), bounds_t2, ind2)
-
-         CALL mp_sum(ind1, mp_comm)
-         CALL mp_sum(ind2, mp_comm)
-
-         IF (.NOT. ANY(ind1 .AND. ind2)) THEN
-            dbcsr_t_need_contract = .FALSE.
-            CALL timestop(handle)
-            RETURN
-         ENDIF
-
-         DEALLOCATE (ind1, ind2)
-      ENDDO
-
-      CALL timestop(handle)
-   END FUNCTION
-
    SUBROUTINE dbcsr_t_map_bounds_to_tensors(tensor_1, tensor_2, &
                                             contract_1, notcontract_1, &
                                             contract_2, notcontract_2, &

src/tensors/dbcsr_tensor_api.F

@@ -18,7 +18,7 @@ MODULE dbcsr_tensor_api
    USE dbcsr_tensor, ONLY: &
       dbcsr_t_contract, dbcsr_t_get_block, dbcsr_t_get_stored_coordinates, dbcsr_t_put_block, &
       dbcsr_t_reserve_blocks, dbcsr_t_copy_matrix_to_tensor, dbcsr_t_copy, &
-      dbcsr_t_copy_tensor_to_matrix, dbcsr_t_need_contract, dbcsr_t_batched_contract_init, &
+      dbcsr_t_copy_tensor_to_matrix, dbcsr_t_batched_contract_init, &
       dbcsr_t_batched_contract_finalize, &
       dbcsr_t_contract_index
    USE dbcsr_tensor_block, ONLY: &
@@ -84,7 +84,6 @@ MODULE dbcsr_tensor_api
    PUBLIC :: dbcsr_t_get_info
    PUBLIC :: dbcsr_t_checksum
    PUBLIC :: dbcsr_t_finalize
-   PUBLIC :: dbcsr_t_need_contract
    PUBLIC :: dbcsr_t_scale
    PUBLIC :: dbcsr_t_get_num_blocks, dbcsr_t_get_num_blocks_total
    PUBLIC :: dbcsr_t_get_nze, dbcsr_t_get_nze_total

src/work/dbcsr_work_operations.F

@@ -288,9 +288,8 @@ SUBROUTINE dbcsr_create_new(matrix, name, dist, matrix_type, &
       ELSEIF (PRESENT(col_blk_size)) THEN
          CALL array_new(matrix%col_blk_size, col_blk_size, gift=reuse_arrays)
       ELSE
-         DBCSR_ABORT("Missing row_blk_size")
+         DBCSR_ABORT("Missing col_blk_size")
       ENDIF
-
       IF (PRESENT(max_cbs)) THEN
          matrix%max_cbs = max_cbs
       ELSE IF (array_size(matrix%col_blk_size) .GT. 0) THEN
@@ -490,15 +489,20 @@ SUBROUTINE dbcsr_create_template(matrix, template, name, dist, matrix_type, &
       ELSE
          new_matrix_type = dbcsr_get_matrix_type(template)
       ENDIF
+      !
+      IF ((PRESENT(row_blk_size) .NEQV. PRESENT(col_blk_size)) .OR. &
+          (PRESENT(row_blk_size_obj) .NEQV. PRESENT(col_blk_size_obj))) THEN
+         DBCSR_ABORT("Both row_blk_size and col_blk_size must be provided")
+      ENDIF
+      !
       IF (PRESENT(max_rbs)) new_max_rbs = max_rbs
       IF (PRESENT(row_blk_offset)) new_row_blk_offset = row_blk_offset
       NULLIFY (blk_size)
-      IF (PRESENT(row_blk_size)) THEN
-         blk_size => row_blk_size
-      ELSEIF (PRESENT(row_blk_size_obj)) THEN
+      IF (PRESENT(row_blk_size_obj)) THEN
          blk_size => array_data(row_blk_size_obj)
+      ELSEIF (PRESENT(row_blk_size)) THEN
+         blk_size => row_blk_size
       ENDIF
-
       IF (ASSOCIATED(blk_size)) THEN
          IF (.NOT. PRESENT(max_rbs)) &
             new_max_rbs = MAXVAL(blk_size)
@@ -512,10 +516,10 @@ SUBROUTINE dbcsr_create_template(matrix, template, name, dist, matrix_type, &
       IF (PRESENT(max_cbs)) new_max_cbs = max_cbs
       IF (PRESENT(col_blk_offset)) new_col_blk_offset = col_blk_offset
       NULLIFY (blk_size)
-      IF (PRESENT(col_blk_size)) THEN
-         blk_size => col_blk_size
-      ELSEIF (PRESENT(col_blk_size_obj)) THEN
+      IF (PRESENT(col_blk_size_obj)) THEN
          blk_size => array_data(col_blk_size_obj)
+      ELSEIF (PRESENT(col_blk_size)) THEN
+         blk_size => col_blk_size
       ENDIF
       IF (ASSOCIATED(blk_size)) THEN
          IF (.NOT. PRESENT(max_cbs)) &
@@ -551,11 +555,11 @@ SUBROUTINE dbcsr_create_template(matrix, template, name, dist, matrix_type, &
       ELSE
          new_mutable_work = dbcsr_use_mutable(template)
       ENDIF
-      IF (PRESENT(row_blk_size) .AND. PRESENT(col_blk_size)) THEN
+      IF (PRESENT(row_blk_size_obj)) THEN
          CALL dbcsr_create(matrix, name=new_name, dist=new_dist, &
                            matrix_type=new_matrix_type, &
-                           row_blk_size=row_blk_size, &
-                           col_blk_size=col_blk_size, &
+                           row_blk_size_obj=row_blk_size_obj, &
+                           col_blk_size_obj=col_blk_size_obj, &
                            nze=nze, &
                            data_type=new_data_type, &
                            data_buffer=data_buffer, &
@@ -567,11 +571,11 @@ SUBROUTINE dbcsr_create_template(matrix, template, name, dist, matrix_type, &
                            mutable_work=new_mutable_work, &
                            make_index=make_index, &
                            replication_type=new_replication_type)
-      ELSEIF (PRESENT(row_blk_size_obj) .AND. PRESENT(col_blk_size_obj)) THEN
+      ELSEIF (PRESENT(row_blk_size)) THEN
          CALL dbcsr_create(matrix, name=new_name, dist=new_dist, &
                            matrix_type=new_matrix_type, &
-                           row_blk_size_obj=row_blk_size_obj, &
-                           col_blk_size_obj=col_blk_size_obj, &
+                           row_blk_size=row_blk_size, &
+                           col_blk_size=col_blk_size, &
                            nze=nze, &
                            data_type=new_data_type, &
                            data_buffer=data_buffer, &