Address comments

codeplaysoftware · Aug 2, 2023 · 831ab63 · 831ab63
1 parent 377fa8d
commit 831ab63
Show file tree

Hide file tree

Showing 59 changed files with 460 additions and 1,236 deletions.
diff --git a/cmake/CmakeFunctionHelper.cmake b/cmake/CmakeFunctionHelper.cmake
@@ -109,7 +109,6 @@ function(generate_blas_unary_objects blas_level func)
                 ${cpp_data}
                 ${index}
                 ${increment}
-                ${cpp_data}
                 ${file_name}
                 MAIN_DEPENDENCY ${SYCLBLAS_SRC}/interface/${blas_level}/${func}.cpp.in
                 DEPENDS ${SYCLBLAS_SRC_GENERATOR}/py_gen_blas_unary.py
@@ -128,47 +127,6 @@ function(generate_blas_unary_objects blas_level func)
   add_sycl_to_target(TARGET ${func} SOURCES ${FUNC_SRC})
 endfunction(generate_blas_unary_objects)
 
-# blas binary function for generating source code
-function(generate_blas_binary_objects blas_level func)
-  set(LOCATION "${SYCLBLAS_GENERATED_SRC}/${blas_level}/${func}/")
-  foreach(data ${data_list})
-    cpp_type(cpp_data ${data})
-    foreach(index ${index_list})
-      set(container_names "${data}_${data}")
-      foreach(increment ${index_list})
-        sanitize_file_name(file_name
-                "${func}_${data}_${index}_${container_names}_${increment}.cpp")
-        add_custom_command(OUTPUT "${LOCATION}/${file_name}"
-                COMMAND ${PYTHON_EXECUTABLE} ${SYCLBLAS_SRC_GENERATOR}/py_gen_blas_binary.py
-                ${PROJECT_SOURCE_DIR}/external/
-                ${SYCLBLAS_SRC_GENERATOR}/gen
-                ${blas_level}
-                ${func}
-                ${SYCLBLAS_SRC}/interface/${blas_level}/${func}.cpp.in
-                ${cpp_data}
-                ${index}
-                ${increment}
-                ${cpp_data}
-                ${cpp_data}
-                ${file_name}
-                MAIN_DEPENDENCY ${SYCLBLAS_SRC}/interface/${blas_level}/${func}.cpp.in
-                DEPENDS ${SYCLBLAS_SRC_GENERATOR}/py_gen_blas_binary.py
-                WORKING_DIRECTORY ${PROJECT_BINARY_DIR}
-                VERBATIM
-                )
-        list(APPEND FUNC_SRC "${LOCATION}/${file_name}")
-      endforeach(increment)
-    endforeach(index)
-  endforeach(data)
-  add_library(${func} OBJECT ${FUNC_SRC})
-  set_target_compile_def(${func})
-  target_include_directories(${func} PRIVATE ${SYCLBLAS_SRC} ${SYCLBLAS_INCLUDE}
-          ${SYCLBLAS_COMMON_INCLUDE_DIR} ${THIRD_PARTIES_INCLUDE})
-  message(STATUS "Adding SYCL to target ${func}")
-  add_sycl_to_target(TARGET ${func} SOURCES ${FUNC_SRC})
-endfunction(generate_blas_binary_objects)
-
-
 # blas binary function for generating source code
 function(generate_blas_reduction_objects blas_level func)
 set(LOCATION "${SYCLBLAS_GENERATED_SRC}/${blas_level}/${func}/")
@@ -190,8 +148,6 @@ foreach(data ${data_list})
                 ${cpp_data}
                 ${index}
                 ${increment}
-                ${cpp_data}
-                ${cpp_data}
                 ${operator}
                 ${file_name}
                 MAIN_DEPENDENCY ${SYCLBLAS_SRC}/interface/${blas_level}/${func}.cpp.in
@@ -212,88 +168,6 @@ message(STATUS "Adding SYCL to target ${func}")
 add_sycl_to_target(TARGET ${func} SOURCES ${FUNC_SRC})
 endfunction(generate_blas_reduction_objects)
 
-# blas special binary function for generating source code
-function(generate_blas_binary_special_objects blas_level func)
-  set(LOCATION "${SYCLBLAS_GENERATED_SRC}/${blas_level}/${func}/")
-  foreach(data ${data_list})
-    cpp_type(cpp_data ${data})
-    foreach(index ${index_list})
-      set(container_names "${data}_${data}")
-      foreach(increment ${index_list})
-        sanitize_file_name(file_name
-                "${func}_${data}_${index}_${container_names}_${increment}.cpp")
-        add_custom_command(OUTPUT "${LOCATION}/${file_name}"
-                COMMAND ${PYTHON_EXECUTABLE} ${SYCLBLAS_SRC_GENERATOR}/py_gen_blas_binary_special.py
-                ${PROJECT_SOURCE_DIR}/external/
-                ${SYCLBLAS_SRC_GENERATOR}/gen
-                ${blas_level}
-                ${func}
-                ${SYCLBLAS_SRC}/interface/${blas_level}/${func}.cpp.in
-                ${cpp_data}
-                ${index}
-                ${increment}
-                ${cpp_data}
-                ${cpp_data}
-                ${file_name}
-                MAIN_DEPENDENCY ${SYCLBLAS_SRC}/interface/${blas_level}/${func}.cpp.in
-                DEPENDS ${SYCLBLAS_SRC_GENERATOR}/py_gen_blas_binary_special.py
-                WORKING_DIRECTORY ${PROJECT_BINARY_DIR}
-                VERBATIM
-                )
-        list(APPEND FUNC_SRC "${LOCATION}/${file_name}")
-      endforeach(increment)
-    endforeach(index)
-  endforeach(data)
-  add_library(${func} OBJECT ${FUNC_SRC})
-  set_target_compile_def(${func})
-  target_include_directories(${func} PRIVATE ${SYCLBLAS_SRC} ${SYCLBLAS_INCLUDE}
-          ${SYCLBLAS_COMMON_INCLUDE_DIR} ${THIRD_PARTIES_INCLUDE})
-  message(STATUS "Adding SYCL to target ${func}")
-  add_sycl_to_target(TARGET ${func} SOURCES ${FUNC_SRC})
-endfunction(generate_blas_binary_special_objects)
-
-# blas ternary function for generating source code
-function(generate_blas_ternary_objects blas_level func)
-  set(LOCATION "${SYCLBLAS_GENERATED_SRC}/${blas_level}/${func}/")
-  foreach(data ${data_list})
-    cpp_type(cpp_data ${data})
-    foreach(index ${index_list})
-      set(container_names
-              "${data}_${data}_${container2}")
-      foreach(increment ${index_list})
-        sanitize_file_name(file_name
-                "${func}_${data}_${index}_${container_names}_${increment}.cpp")
-        add_custom_command(OUTPUT "${LOCATION}/${file_name}"
-                COMMAND ${PYTHON_EXECUTABLE} ${SYCLBLAS_SRC_GENERATOR}/py_gen_blas_ternary.py
-                ${PROJECT_SOURCE_DIR}/external/
-                ${SYCLBLAS_SRC_GENERATOR}/gen
-                ${blas_level}
-                ${func}
-                ${SYCLBLAS_SRC}/interface/${blas_level}/${func}.cpp.in
-                ${cpp_data}
-                ${index}
-                ${increment}
-                ${cpp_data}
-                ${cpp_data}
-                ${cpp_data}
-                ${file_name}
-                MAIN_DEPENDENCY ${SYCLBLAS_SRC}/interface/${blas_level}/${func}.cpp.in
-                DEPENDS ${SYCLBLAS_SRC_GENERATOR}/py_gen_blas_ternary.py
-                WORKING_DIRECTORY ${PROJECT_BINARY_DIR}
-                VERBATIM
-                )
-        list(APPEND FUNC_SRC "${LOCATION}/${file_name}")
-      endforeach(increment)
-    endforeach(index)
-  endforeach(data)
-  add_library(${func} OBJECT ${FUNC_SRC})
-  set_target_compile_def(${func})
-  target_include_directories(${func} PRIVATE ${SYCLBLAS_SRC} ${SYCLBLAS_INCLUDE}
-          ${SYCLBLAS_COMMON_INCLUDE_DIR} ${THIRD_PARTIES_INCLUDE})
-  message(STATUS "Adding SYCL to target ${func}")
-  add_sycl_to_target(TARGET ${func} SOURCES ${FUNC_SRC})
-endfunction(generate_blas_ternary_objects)
-
 # blas function for generating source code for the rotg operator (asynchronous version with containers)
 function(generate_blas_rotg_objects blas_level func)
   set(LOCATION "${SYCLBLAS_GENERATED_SRC}/${blas_level}/${func}/")
@@ -310,10 +184,6 @@ function(generate_blas_rotg_objects blas_level func)
             ${func}
             ${SYCLBLAS_SRC}/interface/${blas_level}/${func}.cpp.in
             ${cpp_data}
-            ${cpp_data}
-            ${cpp_data}
-            ${cpp_data}
-            ${cpp_data}
             ${file_name}
             MAIN_DEPENDENCY ${SYCLBLAS_SRC}/interface/${blas_level}/${func}.cpp.in
             DEPENDS ${SYCLBLAS_SRC_GENERATOR}/py_gen_blas_rotg.py
@@ -330,73 +200,6 @@ function(generate_blas_rotg_objects blas_level func)
   add_sycl_to_target(TARGET ${func} SOURCES ${FUNC_SRC})
 endfunction(generate_blas_rotg_objects)
 
-# blas function for generating source code for the rotg operator (synchronous version)
-function(generate_blas_rotg_return_objects blas_level func)
-  set(LOCATION "${SYCLBLAS_GENERATED_SRC}/${blas_level}/${func}/")
-  foreach (data ${data_list})
-    cpp_type(cpp_data ${data})
-    sanitize_file_name(file_name
-            "${func}_${data}_${index}_${data}_${increment}.cpp")
-    add_custom_command(OUTPUT "${LOCATION}/${file_name}"
-            COMMAND ${PYTHON_EXECUTABLE} ${SYCLBLAS_SRC_GENERATOR}/py_gen_blas_rotg_return.py
-            ${PROJECT_SOURCE_DIR}/external/
-            ${SYCLBLAS_SRC_GENERATOR}/gen
-            ${blas_level}
-            ${func}
-            ${SYCLBLAS_SRC}/interface/${blas_level}/${func}.cpp.in
-            ${cpp_data}
-            ${file_name}
-            MAIN_DEPENDENCY ${SYCLBLAS_SRC}/interface/${blas_level}/${func}.cpp.in
-            DEPENDS ${SYCLBLAS_SRC_GENERATOR}/py_gen_blas_rotg_return.py
-            WORKING_DIRECTORY ${PROJECT_BINARY_DIR}
-            VERBATIM
-            )
-    list(APPEND FUNC_SRC "${LOCATION}/${file_name}")
-  endforeach (data)
-  add_library(${func} OBJECT ${FUNC_SRC})
-  set_target_compile_def(${func})
-  target_include_directories(${func} PRIVATE ${SYCLBLAS_SRC} ${SYCLBLAS_INCLUDE}
-          ${SYCLBLAS_COMMON_INCLUDE_DIR} ${THIRD_PARTIES_INCLUDE})
-  message(STATUS "Adding SYCL to target ${func}")
-  add_sycl_to_target(TARGET ${func} SOURCES ${FUNC_SRC})
-endfunction(generate_blas_rotg_return_objects)
-
-# blas function for generating source code for the rotg operator (asynchronous version with containers)
-function(generate_blas_rotmg_objects blas_level func)
-  set(LOCATION "${SYCLBLAS_GENERATED_SRC}/${blas_level}/${func}/")
-  foreach (data ${data_list})
-    cpp_type(cpp_data ${data})
-    set(container_names "${data}_${data}_${data}_${data}_${data}")
-    sanitize_file_name(file_name "${func}_${data}_${container_names}.cpp")
-    add_custom_command(OUTPUT "${LOCATION}/${file_name}"
-            COMMAND ${PYTHON_EXECUTABLE} ${SYCLBLAS_SRC_GENERATOR}/py_gen_blas_rotmg.py
-            ${PROJECT_SOURCE_DIR}/external/
-            ${SYCLBLAS_SRC_GENERATOR}/gen
-            ${blas_level}
-            ${func}
-            ${SYCLBLAS_SRC}/interface/${blas_level}/${func}.cpp.in
-            ${cpp_data}
-            ${data}
-            ${data}
-            ${data}
-            ${data}
-            ${data}
-            ${file_name}
-            MAIN_DEPENDENCY ${SYCLBLAS_SRC}/interface/${blas_level}/${func}.cpp.in
-            DEPENDS ${SYCLBLAS_SRC_GENERATOR}/py_gen_blas_rotmg.py
-            WORKING_DIRECTORY ${PROJECT_BINARY_DIR}
-            VERBATIM
-            )
-    list(APPEND FUNC_SRC "${LOCATION}/${file_name}")
-  endforeach (data)
-  add_library(${func} OBJECT ${FUNC_SRC})
-  set_target_compile_def(${func})
-  target_include_directories(${func} PRIVATE ${SYCLBLAS_SRC} ${SYCLBLAS_INCLUDE}
-          ${SYCLBLAS_COMMON_INCLUDE_DIR} ${THIRD_PARTIES_INCLUDE})
-  message(STATUS "Adding SYCL to target ${func}")
-  add_sycl_to_target(TARGET ${func} SOURCES ${FUNC_SRC})
-endfunction(generate_blas_rotmg_objects)
-
 # blas gemm function for generating source code
 function(generate_blas_gemm_objects blas_level func)
 set(LOCATION "${SYCLBLAS_GENERATED_SRC}/${blas_level}/${func}/")
@@ -498,9 +301,6 @@ function(add_gemm_configuration
                       ${use_joint_matrix}
                       ${symm_a}
                       ${symm_b}
-                      ${cpp_data}
-                      ${cpp_data}
-                      ${cpp_data}
                       MAIN_DEPENDENCY ${SYCLBLAS_SRC}/interface/${blas_level}/${func}.cpp.in
                       DEPENDS ${SYCLBLAS_SRC_GENERATOR}/py_gen_blas_gemm_launcher.py
                       WORKING_DIRECTORY ${PROJECT_BINARY_DIR}

diff --git a/doc/AddingBlas3Op.md b/doc/AddingBlas3Op.md
@@ -197,28 +197,24 @@ namespace internal {
 template typename SB_Handle::event_t _trsm(
   SB_Handle sb_handle, char Side, char Triangle, char Transpose, char Diagonal,
   ${INDEX_TYPE} M, ${INDEX_TYPE} N, ${DATA_TYPE} alpha,
-  ${container_t0} A, ${INDEX_TYPE} lda,
-  ${container_t1} B, ${INDEX_TYPE} ldb);
+  BufferIterator<${DATA_TYPE}> A, ${INDEX_TYPE} lda,
+  BufferIterator<${DATA_TYPE}> B, ${INDEX_TYPE} ldb);
 
 
 } // namespace internal
 } // namespace blas
 ```
 
-Where `${INDEX_TYPE}, ${DATA_TYPE}, ${container_t0}` and `${container_t1}` are going
+Where `${INDEX_TYPE} and ${DATA_TYPE} are going
 to be replaced by the appropriate types required to explicitly instantiate the new function
 Finally, the file `src/interface/blas3/CMakeLists.txt` must be changed
 in order to generate instantiations of `_trsm`.
 The following entry must be added:
 
 ```cmake
-generate_blas_binary_objects(blas3 trsm)
+generate_blas_unary_objects(blas3 trsm)
 ```
 
-There are predefined functions to be used depending on the number of inputs the function expects.
-`_trsm` is a case of *binary* function, since only two buffers are required as input. `_gemm` is
-an example of ternary function, since it requires three buffers as input (in this case `${container_t2}`
-can be used in the `.cpp` file).
 
 After this, the new object file created must be added to the library and is done by adding a new
 entry to `cmake/CmakeFunctionHelper.cmake`. At the end of this file there is a list of all object

diff --git a/doc/Gemm.md b/doc/Gemm.md
@@ -249,15 +249,15 @@ namespace internal {
 // gemm
 template typename SB_Handle::event_t _gemm(
     SB_Handle& sb_handle, char _TransA, char _TransB, ${INDEX_TYPE} _M,
-    ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, ${DATA_TYPE} _alpha, ${container_t0} a_,
-    ${INDEX_TYPE} _lda, ${container_t1} b_, ${INDEX_TYPE} _ldb,
-    ${DATA_TYPE} _beta, ${container_t2} _C, ${INDEX_TYPE} _ldc);
+    ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, ${DATA_TYPE} _alpha, ${DATA_TYPE} a_,
+    ${INDEX_TYPE} _lda, ${DATA_TYPE} b_, ${INDEX_TYPE} _ldb,
+    ${DATA_TYPE} _beta, ${DATA_TYPE} _C, ${INDEX_TYPE} _ldc);
 // batched gemm
 template typename SB_Handle::event_t _gemm_batched(
     SB_Handle& sb_handle, char _TransA, char _TransB, ${INDEX_TYPE} _M,
-    ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, ${DATA_TYPE} _alpha, ${container_t0} a_,
-    ${INDEX_TYPE} _lda, ${INDEX_TYPE} _stridea, ${container_t1} b_, ${INDEX_TYPE} _ldb, 
-    ${INDEX_TYPE} _strideb, ${DATA_TYPE} _beta, ${container_t2} _C, ${INDEX_TYPE} _ldc,
+    ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, ${DATA_TYPE} _alpha, ${DATA_TYPE} a_,
+    ${INDEX_TYPE} _lda, ${INDEX_TYPE} _stridea, ${DATA_TYPE} b_, ${INDEX_TYPE} _ldb, 
+    ${INDEX_TYPE} _strideb, ${DATA_TYPE} _beta, ${DATA_TYPE} _C, ${INDEX_TYPE} _ldc,
     ${INDEX_TYPE} _stridec, ${INDEX_TYPE} batch_size, gemm_batch_type_t batch_type);
 }  // namespace internal
 }  // namespace blas

diff --git a/include/blas_meta.h b/include/blas_meta.h
@@ -111,19 +111,6 @@ struct RebindType {
   using type = RemoveAll<element_t> *;
 };
 
-template <typename element_t, bool isConst>
-struct ReturnType;
-
-template <typename element_t>
-struct ReturnType <element_t, true> {
-  using type = const element_t;
-};
-
-template <typename element_t>
-struct ReturnType <element_t, false> {
-  using type = element_t;
-};
-
 template <typename index_t>
 inline bool is_power_of_2(index_t ind) {
   return ind > 0 && !(ind & (ind - 1));

diff --git a/include/operations/blas2_trees.h b/include/operations/blas2_trees.h
@@ -116,7 +116,7 @@ struct SumMatrixColumns {
 
   bool valid_thread(cl::sycl::nd_item<1> ndItem) const;
 
-  value_t eval(index_t i) const;
+  value_t eval(index_t i);
 
   value_t eval(cl::sycl::nd_item<1> ndItem) const;
   void bind(cl::sycl::handler &h);