From 831ab631298d832784102db06bbe9ffa874859f5 Mon Sep 17 00:00:00 2001 From: Alejandro Acosta Date: Wed, 2 Aug 2023 13:53:03 +0100 Subject: [PATCH] Address comments --- cmake/CmakeFunctionHelper.cmake | 200 ------------------ doc/AddingBlas3Op.md | 12 +- doc/Gemm.md | 12 +- include/blas_meta.h | 13 -- include/operations/blas2_trees.h | 2 +- python_generator/py_gen_blas_binary.py | 94 -------- .../py_gen_blas_binary_special.py | 93 -------- python_generator/py_gen_blas_gemm_launcher.py | 18 -- python_generator/py_gen_blas_reduction.py | 16 +- python_generator/py_gen_blas_rotg.py | 26 +-- python_generator/py_gen_blas_rotg_return.py | 69 ------ python_generator/py_gen_blas_rotmg.py | 99 --------- python_generator/py_gen_blas_ternary.py | 99 --------- python_generator/py_gen_blas_unary.py | 8 +- src/interface/blas1/CMakeLists.txt | 33 ++- src/interface/blas1/asum.cpp.in | 30 ++- src/interface/blas1/asum_return.cpp.in | 12 +- src/interface/blas1/axpy.cpp.in | 19 +- src/interface/blas1/copy.cpp.in | 14 +- src/interface/blas1/dot.cpp.in | 23 +- src/interface/blas1/dot_return.cpp.in | 14 +- src/interface/blas1/iamax.cpp.in | 33 ++- src/interface/blas1/iamax_return.cpp.in | 14 +- src/interface/blas1/iamin.cpp.in | 33 ++- src/interface/blas1/iamin_return.cpp.in | 14 +- src/interface/blas1/nrm2.cpp.in | 30 ++- src/interface/blas1/nrm2_return.cpp.in | 12 +- src/interface/blas1/rot.cpp.in | 24 +-- src/interface/blas1/rotg.cpp.in | 21 +- src/interface/blas1/rotg_return.cpp.in | 5 +- src/interface/blas1/rotm.cpp.in | 20 +- src/interface/blas1/rotmg.cpp.in | 9 +- src/interface/blas1/scal.cpp.in | 20 +- src/interface/blas1/sdsdot.cpp.in | 14 +- src/interface/blas1/sdsdot_return.cpp.in | 13 +- src/interface/blas1/swap.cpp.in | 10 +- src/interface/blas2/CMakeLists.txt | 30 +-- src/interface/blas2/gbmv.cpp.in | 16 +- src/interface/blas2/gemv.cpp.in | 26 +-- src/interface/blas2/ger.cpp.in | 13 +- src/interface/blas2/sbmv.cpp.in | 19 +- src/interface/blas2/spmv.cpp.in | 15 +- src/interface/blas2/spr.cpp.in | 9 +- src/interface/blas2/spr2.cpp.in | 15 +- src/interface/blas2/symv.cpp.in | 16 +- src/interface/blas2/syr.cpp.in | 9 +- src/interface/blas2/syr2.cpp.in | 17 +- src/interface/blas2/tbmv.cpp.in | 12 +- src/interface/blas2/tbsv.cpp.in | 12 +- src/interface/blas2/tpmv.cpp.in | 8 +- src/interface/blas2/trmv.cpp.in | 9 +- src/interface/blas2/trsv.cpp.in | 9 +- src/interface/blas3/CMakeLists.txt | 6 +- src/interface/blas3/gemm.cpp.in | 103 +++++---- src/interface/blas3/gemm_launcher.cpp.in | 107 +++------- src/interface/blas3/symm.cpp.in | 18 +- src/interface/blas3/trsm.cpp.in | 15 +- src/interface/reduction/reduction.cpp.in | 32 +-- src/operations/blas2/gemv.hpp | 2 +- 59 files changed, 460 insertions(+), 1236 deletions(-) delete mode 100644 python_generator/py_gen_blas_binary.py delete mode 100644 python_generator/py_gen_blas_binary_special.py delete mode 100644 python_generator/py_gen_blas_rotg_return.py delete mode 100644 python_generator/py_gen_blas_rotmg.py delete mode 100644 python_generator/py_gen_blas_ternary.py diff --git a/cmake/CmakeFunctionHelper.cmake b/cmake/CmakeFunctionHelper.cmake index 30f1d0185..67866b2e0 100644 --- 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 index 11867fc88..b7283c4ce 100644 --- 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 index 4424cd2cc..b69aabe14 100644 --- 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 index fd2debb4d..93c98c6e1 100644 --- a/include/blas_meta.h +++ b/include/blas_meta.h @@ -111,19 +111,6 @@ struct RebindType { using type = RemoveAll *; }; -template -struct ReturnType; - -template -struct ReturnType { - using type = const element_t; -}; - -template -struct ReturnType { - using type = element_t; -}; - template 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 index d86ca61b3..8f7f8c678 100644 --- 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); diff --git a/python_generator/py_gen_blas_binary.py b/python_generator/py_gen_blas_binary.py deleted file mode 100644 index 0fbb383b8..000000000 --- a/python_generator/py_gen_blas_binary.py +++ /dev/null @@ -1,94 +0,0 @@ -#/*************************************************************************** -# * -# * @license -# * Copyright (C) Codeplay Software Limited -# * Licensed under the Apache License, Version 2.0 (the "License"); -# * you may not use this file except in compliance with the License. -# * You may obtain a copy of the License at -# * -# * http://www.apache.org/licenses/LICENSE-2.0 -# * -# * For your convenience, a copy of the License has been included in this -# * repository. -# * -# * Unless required by applicable law or agreed to in writing, software -# * distributed under the License is distributed on an "AS IS" BASIS, -# * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -# * See the License for the specific language governing permissions and -# * limitations under the License. -# * -# * SYCL-BLAS: BLAS implementation using SYCL -# * -# * @filename py_gen_blas_binary.py -# * -# **************************************************************************/ - -# py_gen import -import errno -import os -import sys - -if __name__ == '__main__': - - generator_path = sys.argv[1] - sys.path.insert(0, generator_path) - from py_gen import generate_file - from py_gen import * - from string import Template - - index_list = ['int', 'long', 'long long'] - - input_template = sys.argv[2] - blas_level_name = sys.argv[3] - blas_function_name = sys.argv[4] - blas_template_impl = sys.argv[5] - data = sys.argv[6] - index = sys.argv[7] - increment = sys.argv[8] - container0 = sys.argv[9] - container1 = sys.argv[10] - file_name = sys.argv[11] - source = 'generated_src/' + blas_level_name + '/' + blas_function_name + '/' - - try: - os.makedirs(source) - except OSError as e: - if e.errno != errno.EEXIST: - raise - f = open(blas_template_impl, "r") - template = Template(f.read()) - f.close() - iterables = [ - Iterable( - key='DATA_TYPE', - vals=[data], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='INDEX_TYPE', - vals=[index], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='INCREMENT_TYPE', - vals=[increment], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t0', - vals=[container0], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t1', - vals=[container1], - itermode=Itermode.combinations, - iter_modifier=1) - ] - iter_groups = [IterGroup('@ip1@', template, iterables, combine_iters=True)] - generate_file( - input_template, - source + file_name, - iter_groups, - format_generated=False, - format_script="") diff --git a/python_generator/py_gen_blas_binary_special.py b/python_generator/py_gen_blas_binary_special.py deleted file mode 100644 index 8d32cf7e9..000000000 --- a/python_generator/py_gen_blas_binary_special.py +++ /dev/null @@ -1,93 +0,0 @@ -#/*************************************************************************** -# * -# * @license -# * Copyright (C) Codeplay Software Limited -# * Licensed under the Apache License, Version 2.0 (the "License"); -# * you may not use this file except in compliance with the License. -# * You may obtain a copy of the License at -# * -# * http://www.apache.org/licenses/LICENSE-2.0 -# * -# * For your convenience, a copy of the License has been included in this -# * repository. -# * -# * Unless required by applicable law or agreed to in writing, software -# * distributed under the License is distributed on an "AS IS" BASIS, -# * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -# * See the License for the specific language governing permissions and -# * limitations under the License. -# * -# * SYCL-BLAS: BLAS implementation using SYCL -# * -# * @filename py_gen_blas_binary_special.py -# * -# **************************************************************************/ -# py_gen import -import errno -import os -import sys - -if __name__ == '__main__': - - generator_path = sys.argv[1] - sys.path.insert(0, generator_path) - from py_gen import generate_file - from py_gen import * - from string import Template - - index_list = ['int', 'long', 'long long'] - - input_template = sys.argv[2] - blas_level_name = sys.argv[3] - blas_function_name = sys.argv[4] - blas_template_impl = sys.argv[5] - data = sys.argv[6] - index = sys.argv[7] - increment = sys.argv[8] - container0 = sys.argv[9] - container1 = sys.argv[10] - file_name = sys.argv[11] - source = 'generated_src/' + blas_level_name + '/' + blas_function_name + '/' - - try: - os.makedirs(source) - except OSError as e: - if e.errno != errno.EEXIST: - raise - f = open(blas_template_impl, "r") - template = Template(f.read()) - f.close() - iterables = [ - Iterable( - key='DATA_TYPE', - vals=[data], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='INDEX_TYPE', - vals=[index], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='INCREMENT_TYPE', - vals=[increment], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t0', - vals=[container0], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t1', - vals=[container1], - itermode=Itermode.combinations, - iter_modifier=1) - ] - iter_groups = [IterGroup('@ip1@', template, iterables, combine_iters=True)] - generate_file( - input_template, - source + file_name, - iter_groups, - format_generated=False, - format_script="") diff --git a/python_generator/py_gen_blas_gemm_launcher.py b/python_generator/py_gen_blas_gemm_launcher.py index b77a6e997..9bac81c4d 100644 --- a/python_generator/py_gen_blas_gemm_launcher.py +++ b/python_generator/py_gen_blas_gemm_launcher.py @@ -72,9 +72,6 @@ use_joint_matrix = sys.argv[37] symm_a = sys.argv[38] symm_b = sys.argv[39] - container0 = sys.argv[40] - container1 = sys.argv[41] - container2 = sys.argv[42] source = 'generated_src/' + blas_level_name + '/' + blas_function_name + '/' try: os.makedirs(source) @@ -249,21 +246,6 @@ key='SYMM_B', vals=[symm_b], itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t0', - vals=[container0], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t1', - vals=[container1], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t2', - vals=[container2], - itermode=Itermode.combinations, iter_modifier=1) ] iter_groups = [IterGroup('@ip1@', template, iterables, combine_iters=True)] diff --git a/python_generator/py_gen_blas_reduction.py b/python_generator/py_gen_blas_reduction.py index d9f43d23a..93be4c76a 100644 --- a/python_generator/py_gen_blas_reduction.py +++ b/python_generator/py_gen_blas_reduction.py @@ -45,10 +45,8 @@ data = sys.argv[6] index = sys.argv[7] increment = sys.argv[8] - container0 = sys.argv[9] - container1 = sys.argv[10] - operator = sys.argv[11] - file_name = sys.argv[12] + operator = sys.argv[9] + file_name = sys.argv[10] source = 'generated_src/' + blas_level_name + '/' + blas_function_name + '/' try: @@ -75,16 +73,6 @@ vals=[increment], itermode=Itermode.combinations, iter_modifier=1), - Iterable( - key='container_t0', - vals=[container0], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t1', - vals=[container1], - itermode=Itermode.combinations, - iter_modifier=1), Iterable( key='OPERATOR', vals=[operator], diff --git a/python_generator/py_gen_blas_rotg.py b/python_generator/py_gen_blas_rotg.py index d441c7709..561832153 100644 --- a/python_generator/py_gen_blas_rotg.py +++ b/python_generator/py_gen_blas_rotg.py @@ -42,11 +42,7 @@ blas_function_name = sys.argv[4] blas_template_impl = sys.argv[5] data = sys.argv[6] - container0 = sys.argv[7] - container1 = sys.argv[8] - container2 = sys.argv[9] - container3= sys.argv[10] - file_name = sys.argv[11] + file_name = sys.argv[7] source = 'generated_src/' + blas_level_name + '/' + blas_function_name + '/' try: @@ -62,26 +58,6 @@ key='DATA_TYPE', vals=[data], itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t0', - vals=[container0], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t1', - vals=[container1], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t2', - vals=[container2], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t3', - vals=[container3], - itermode=Itermode.combinations, iter_modifier=1) ] iter_groups = [IterGroup('@ip1@', template, iterables, combine_iters=True)] diff --git a/python_generator/py_gen_blas_rotg_return.py b/python_generator/py_gen_blas_rotg_return.py deleted file mode 100644 index 7f945c516..000000000 --- a/python_generator/py_gen_blas_rotg_return.py +++ /dev/null @@ -1,69 +0,0 @@ -#/*************************************************************************** -# * -# * @license -# * Copyright (C) Codeplay Software Limited -# * Licensed under the Apache License, Version 2.0 (the "License"); -# * you may not use this file except in compliance with the License. -# * You may obtain a copy of the License at -# * -# * http://www.apache.org/licenses/LICENSE-2.0 -# * -# * For your convenience, a copy of the License has been included in this -# * repository. -# * -# * Unless required by applicable law or agreed to in writing, software -# * distributed under the License is distributed on an "AS IS" BASIS, -# * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -# * See the License for the specific language governing permissions and -# * limitations under the License. -# * -# * SYCL-BLAS: BLAS implementation using SYCL -# * -# * @filename py_gen_blas_rotg_return.py -# * -# **************************************************************************/ -# py_gen import -import errno -import os -import sys - -if __name__ == '__main__': - - generator_path = sys.argv[1] - sys.path.insert(0, generator_path) - from py_gen import generate_file - from py_gen import * - from string import Template - - index_list = ['int', 'long', 'long long'] - - input_template = sys.argv[2] - blas_level_name = sys.argv[3] - blas_function_name = sys.argv[4] - blas_template_impl = sys.argv[5] - data = sys.argv[6] - file_name = sys.argv[7] - source = 'generated_src/' + blas_level_name + '/' + blas_function_name + '/' - - try: - os.makedirs(source) - except OSError as e: - if e.errno != errno.EEXIST: - raise - f = open(blas_template_impl, "r") - template = Template(f.read()) - f.close() - iterables = [ - Iterable( - key='DATA_TYPE', - vals=[data], - itermode=Itermode.combinations, - iter_modifier=1), - ] - iter_groups = [IterGroup('@ip1@', template, iterables, combine_iters=True)] - generate_file( - input_template, - source + file_name, - iter_groups, - format_generated=False, - format_script="") diff --git a/python_generator/py_gen_blas_rotmg.py b/python_generator/py_gen_blas_rotmg.py deleted file mode 100644 index 01bfa16dc..000000000 --- a/python_generator/py_gen_blas_rotmg.py +++ /dev/null @@ -1,99 +0,0 @@ -# /*************************************************************************** -# * -# * @license -# * Copyright (C) Codeplay Software Limited -# * Licensed under the Apache License, Version 2.0 (the "License"); -# * you may not use this file except in compliance with the License. -# * You may obtain a copy of the License at -# * -# * http://www.apache.org/licenses/LICENSE-2.0 -# * -# * For your convenience, a copy of the License has been included in this -# * repository. -# * -# * Unless required by applicable law or agreed to in writing, software -# * distributed under the License is distributed on an "AS IS" BASIS, -# * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -# * See the License for the specific language governing permissions and -# * limitations under the License. -# * -# * SYCL-BLAS: BLAS implementation using SYCL -# * -# * @filename py_gen_blas_rotmg.py -# * -# **************************************************************************/ -# py_gen import -import errno -import os -import sys - -if __name__ == '__main__': - - generator_path = sys.argv[1] - sys.path.insert(0, generator_path) - from py_gen import generate_file - from py_gen import * - from string import Template - - index_list = ['int', 'long', 'long long'] - - input_template = sys.argv[2] - blas_level_name = sys.argv[3] - blas_function_name = sys.argv[4] - blas_template_impl = sys.argv[5] - data = sys.argv[6] - container0 = sys.argv[7] - container1 = sys.argv[8] - container2 = sys.argv[9] - container3 = sys.argv[10] - container4 = sys.argv[11] - file_name = sys.argv[12] - source = 'generated_src/' + blas_level_name + '/' + blas_function_name + '/' - - try: - os.makedirs(source) - except OSError as e: - if e.errno != errno.EEXIST: - raise - f = open(blas_template_impl, "r") - template = Template(f.read()) - f.close() - iterables = [ - Iterable( - key='DATA_TYPE', - vals=[data], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t0', - vals=[container0], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t1', - vals=[container1], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t2', - vals=[container2], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t3', - vals=[container3], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t4', - vals=[container4], - itermode=Itermode.combinations, - iter_modifier=1) - ] - iter_groups = [IterGroup('@ip1@', template, iterables, combine_iters=True)] - generate_file( - input_template, - source + file_name, - iter_groups, - format_generated=False, - format_script="") diff --git a/python_generator/py_gen_blas_ternary.py b/python_generator/py_gen_blas_ternary.py deleted file mode 100644 index 7c44e980a..000000000 --- a/python_generator/py_gen_blas_ternary.py +++ /dev/null @@ -1,99 +0,0 @@ -#/*************************************************************************** -# * -# * @license -# * Copyright (C) Codeplay Software Limited -# * Licensed under the Apache License, Version 2.0 (the "License"); -# * you may not use this file except in compliance with the License. -# * You may obtain a copy of the License at -# * -# * http://www.apache.org/licenses/LICENSE-2.0 -# * -# * For your convenience, a copy of the License has been included in this -# * repository. -# * -# * Unless required by applicable law or agreed to in writing, software -# * distributed under the License is distributed on an "AS IS" BASIS, -# * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -# * See the License for the specific language governing permissions and -# * limitations under the License. -# * -# * SYCL-BLAS: BLAS implementation using SYCL -# * -# * @filename py_gen_blas_ternary.py -# * -# **************************************************************************/ -# py_gen import -import errno -import os -import sys - -if __name__ == '__main__': - - generator_path = sys.argv[1] - sys.path.insert(0, generator_path) - from py_gen import generate_file - from py_gen import * - from string import Template - - index_list = ['int', 'long', 'long long'] - - input_template = sys.argv[2] - blas_level_name = sys.argv[3] - blas_function_name = sys.argv[4] - blas_template_impl = sys.argv[5] - data = sys.argv[6] - index = sys.argv[7] - increment = sys.argv[8] - container0 = sys.argv[9] - container1 = sys.argv[10] - container2 = sys.argv[11] - file_name = sys.argv[12] - source = 'generated_src/' + blas_level_name + '/' + blas_function_name + '/' - - try: - os.makedirs(source) - except OSError as e: - if e.errno != errno.EEXIST: - raise - f = open(blas_template_impl, "r") - template = Template(f.read()) - f.close() - iterables = [ - Iterable( - key='DATA_TYPE', - vals=[data], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='INDEX_TYPE', - vals=[index], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='INCREMENT_TYPE', - vals=[increment], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t0', - vals=[container0], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t1', - vals=[container1], - itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t2', - vals=[container2], - itermode=Itermode.combinations, - iter_modifier=1) - ] - iter_groups = [IterGroup('@ip1@', template, iterables, combine_iters=True)] - generate_file( - input_template, - source + file_name, - iter_groups, - format_generated=False, - format_script="") diff --git a/python_generator/py_gen_blas_unary.py b/python_generator/py_gen_blas_unary.py index deca93835..4286d4b3e 100644 --- a/python_generator/py_gen_blas_unary.py +++ b/python_generator/py_gen_blas_unary.py @@ -44,8 +44,7 @@ data = sys.argv[6] index = sys.argv[7] increment = sys.argv[8] - container0 = sys.argv[9] - file_name = sys.argv[10] + file_name = sys.argv[9] source = 'generated_src/' + blas_level_name + '/' + blas_function_name + '/' try: @@ -71,11 +70,6 @@ key='INCREMENT_TYPE', vals=[increment], itermode=Itermode.combinations, - iter_modifier=1), - Iterable( - key='container_t0', - vals=[container0], - itermode=Itermode.combinations, iter_modifier=1) ] iter_groups = [IterGroup('@ip1@', template, iterables, combine_iters=True)] diff --git a/src/interface/blas1/CMakeLists.txt b/src/interface/blas1/CMakeLists.txt index 1ed769075..a67f9128f 100644 --- a/src/interface/blas1/CMakeLists.txt +++ b/src/interface/blas1/CMakeLists.txt @@ -23,28 +23,27 @@ # * # **************************************************************************/ #blas1 -generate_blas_binary_objects(blas1 axpy) -generate_blas_binary_objects(blas1 asum) -generate_blas_binary_objects(blas1 copy) -generate_blas_binary_objects(blas1 dot_return) -generate_blas_binary_objects(blas1 sdsdot_return) -generate_blas_binary_objects(blas1 nrm2) -generate_blas_binary_objects(blas1 rot) -generate_blas_binary_objects(blas1 nrm2_return) -generate_blas_binary_objects(blas1 swap) +generate_blas_unary_objects(blas1 axpy) +generate_blas_unary_objects(blas1 asum) +generate_blas_unary_objects(blas1 copy) +generate_blas_unary_objects(blas1 dot_return) +generate_blas_unary_objects(blas1 sdsdot_return) +generate_blas_unary_objects(blas1 nrm2) +generate_blas_unary_objects(blas1 rot) +generate_blas_unary_objects(blas1 nrm2_return) +generate_blas_unary_objects(blas1 swap) generate_blas_unary_objects(blas1 asum_return) generate_blas_unary_objects(blas1 iamax_return) generate_blas_unary_objects(blas1 iamin_return) generate_blas_unary_objects(blas1 scal) -generate_blas_ternary_objects(blas1 dot) -generate_blas_ternary_objects(blas1 sdsdot) -generate_blas_ternary_objects(blas1 rotm) -generate_blas_binary_special_objects(blas1 iamax) -generate_blas_binary_special_objects(blas1 iamin) +generate_blas_unary_objects(blas1 dot) +generate_blas_unary_objects(blas1 sdsdot) +generate_blas_unary_objects(blas1 rotm) +generate_blas_unary_objects(blas1 iamax) +generate_blas_unary_objects(blas1 iamin) generate_blas_rotg_objects(blas1 rotg) -generate_blas_rotg_return_objects(blas1 rotg_return) - -generate_blas_rotmg_objects(blas1 rotmg) +generate_blas_rotg_objects(blas1 rotg_return) +generate_blas_rotg_objects(blas1 rotmg) diff --git a/src/interface/blas1/asum.cpp.in b/src/interface/blas1/asum.cpp.in index 0f10b98f8..2aa169f3d 100644 --- a/src/interface/blas1/asum.cpp.in +++ b/src/interface/blas1/asum.cpp.in @@ -38,26 +38,20 @@ namespace internal { * @param _vx VectorView * @param _incx Increment in X axis */ -template typename SB_Handle::event_t _asum(SB_Handle &sb_handle, - ${INDEX_TYPE} _N, - BufferIterator<${container_t0}> _vx, - ${INCREMENT_TYPE} _incx, - BufferIterator<${container_t1}> _rs, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _asum( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, BufferIterator<${DATA_TYPE}> _vx, + ${INCREMENT_TYPE} _incx, BufferIterator<${DATA_TYPE}> _rs, + const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM -template typename SB_Handle::event_t _asum(SB_Handle &sb_handle, - ${INDEX_TYPE} _N, - ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, - ${container_t1}* _rs, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _asum( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _rs, + const typename SB_Handle::event_t& dependencies); -template typename SB_Handle::event_t _asum(SB_Handle &sb_handle, - ${INDEX_TYPE} _N, - const ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, - ${container_t1}* _rs, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _asum( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, const ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _rs, + const typename SB_Handle::event_t& dependencies); #endif } // namespace internal } // namespace blas diff --git a/src/interface/blas1/asum_return.cpp.in b/src/interface/blas1/asum_return.cpp.in index 6cff02525..1b6382cbc 100644 --- a/src/interface/blas1/asum_return.cpp.in +++ b/src/interface/blas1/asum_return.cpp.in @@ -39,17 +39,17 @@ namespace internal { * @param _vx VectorView * @param _incx Increment in X axis */ -template typename ValueType>::type _asum( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, BufferIterator<${container_t0}> _vx, +template typename ValueType>::type _asum( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, BufferIterator<${DATA_TYPE}> _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM -template typename ValueType<${container_t0}*>::type _asum( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, ${container_t0}* _vx, +template typename ValueType<${DATA_TYPE}*>::type _asum( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& dependencies); -template typename ValueType<${container_t0}*>::type _asum( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, const ${container_t0}* _vx, +template typename ValueType<${DATA_TYPE}*>::type _asum( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, const ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& dependencies); #endif diff --git a/src/interface/blas1/axpy.cpp.in b/src/interface/blas1/axpy.cpp.in index 69409e8b3..695e6df89 100644 --- a/src/interface/blas1/axpy.cpp.in +++ b/src/interface/blas1/axpy.cpp.in @@ -39,26 +39,27 @@ namespace internal { * Implements AXPY \f$y = ax + y\f$ * * @param SB_Handle - * @param _vx ${container_t0} + * @param _vx ${DATA_TYPE} * @param _incx Increment in X axis - * @param _vy ${container_t1} + * @param _vy ${DATA_TYPE} * @param _incy Increment in Y axis */ template typename SB_Handle::event_t _axpy( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - BufferIterator<${container_t0}> _vx, ${INCREMENT_TYPE} _incx, BufferIterator<${container_t1}> _vy, - ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& dependencies); + SB_Handle& sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, + BufferIterator<${DATA_TYPE}> _vx, ${INCREMENT_TYPE} _incx, + BufferIterator<${DATA_TYPE}> _vy, ${INCREMENT_TYPE} _incy, + const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _axpy( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - ${container_t0}* _vx, ${INCREMENT_TYPE} _incx, ${container_t1}* _vy, + SB_Handle& sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, + ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& dependencies); template typename SB_Handle::event_t _axpy( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - const ${container_t0}* _vx, ${INCREMENT_TYPE} _incx, ${container_t1}* _vy, + SB_Handle& sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, + const ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& dependencies); #endif diff --git a/src/interface/blas1/copy.cpp.in b/src/interface/blas1/copy.cpp.in index e9762b388..9e912e5c1 100644 --- a/src/interface/blas1/copy.cpp.in +++ b/src/interface/blas1/copy.cpp.in @@ -43,19 +43,19 @@ namespace internal { * @param _incy Increment in Y axis */ template typename SB_Handle::event_t _copy( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, BufferIterator<${container_t0}> _vx, - ${INCREMENT_TYPE} _incx, BufferIterator<${container_t1}> _vy, ${INCREMENT_TYPE} _incy, - const typename SB_Handle::event_t& dependencies); + SB_Handle& sb_handle, ${INDEX_TYPE} _N, BufferIterator<${DATA_TYPE}> _vx, + ${INCREMENT_TYPE} _incx, BufferIterator<${DATA_TYPE}> _vy, + ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _copy( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, ${container_t1}* _vy, ${INCREMENT_TYPE} _incy, + SB_Handle& sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& dependencies); template typename SB_Handle::event_t _copy( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, const ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, ${container_t1}* _vy, ${INCREMENT_TYPE} _incy, + SB_Handle& sb_handle, ${INDEX_TYPE} _N, const ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& dependencies); #endif diff --git a/src/interface/blas1/dot.cpp.in b/src/interface/blas1/dot.cpp.in index 0734747a3..c49e75aa3 100644 --- a/src/interface/blas1/dot.cpp.in +++ b/src/interface/blas1/dot.cpp.in @@ -51,22 +51,17 @@ namespace internal { * @param _rs Output buffer * @return Vector of events to wait for. */ -template typename SB_Handle::event_t _dot(SB_Handle &sb_handle, - ${INDEX_TYPE} _N, BufferIterator<${container_t0}> _vx, - ${INCREMENT_TYPE} _incx, - BufferIterator<${container_t1}> _vy, - ${INCREMENT_TYPE} _incy, - BufferIterator<${container_t2}> _rs, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _dot( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, BufferIterator<${DATA_TYPE}> _vx, + ${INCREMENT_TYPE} _incx, BufferIterator<${DATA_TYPE}> _vy, + ${INCREMENT_TYPE} _incy, BufferIterator<${DATA_TYPE}> _rs, + const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM -template typename SB_Handle::event_t _dot(SB_Handle &sb_handle, - ${INDEX_TYPE} _N, ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, - ${container_t1}* _vy, - ${INCREMENT_TYPE} _incy, - ${container_t2}* _rs, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _dot( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, + ${DATA_TYPE} * _rs, const typename SB_Handle::event_t& dependencies); #endif } // namespace internal diff --git a/src/interface/blas1/dot_return.cpp.in b/src/interface/blas1/dot_return.cpp.in index 36d632b2e..1e51ce345 100644 --- a/src/interface/blas1/dot_return.cpp.in +++ b/src/interface/blas1/dot_return.cpp.in @@ -51,15 +51,15 @@ namespace internal { * @param _rs Output buffer * @return Vector of events to wait for. */ -template typename ValueType>::type _dot( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, BufferIterator<${container_t0}> _vx, - ${INCREMENT_TYPE} _incx, BufferIterator<${container_t1}> _vy, ${INCREMENT_TYPE} _incy, - const typename SB_Handle::event_t& dependencies); +template typename ValueType>::type _dot( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, BufferIterator<${DATA_TYPE}> _vx, + ${INCREMENT_TYPE} _incx, BufferIterator<${DATA_TYPE}> _vy, + ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM -template typename ValueType<${container_t0}>::type _dot( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, ${container_t1}* _vy, ${INCREMENT_TYPE} _incy, +template typename ValueType<${DATA_TYPE}>::type _dot( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& dependencies); #endif diff --git a/src/interface/blas1/iamax.cpp.in b/src/interface/blas1/iamax.cpp.in index ac48660ad..73c885d76 100644 --- a/src/interface/blas1/iamax.cpp.in +++ b/src/interface/blas1/iamax.cpp.in @@ -39,27 +39,24 @@ namespace internal { * @param _vx VectorView * @param _incx Increment in X axis */ -template typename SB_Handle::event_t _iamax(SB_Handle &sb_handle, - ${INDEX_TYPE} _N, - BufferIterator<${container_t0}> _vx, - ${INCREMENT_TYPE} _incx, - BufferIterator> _rs, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _iamax( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, BufferIterator<${DATA_TYPE}> _vx, + ${INCREMENT_TYPE} _incx, + BufferIterator> _rs, + const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM -template typename SB_Handle::event_t _iamax(SB_Handle &sb_handle, - ${INDEX_TYPE} _N, - ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, - IndexValueTuple<${INDEX_TYPE}, ${container_t1}>* _rs, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _iamax( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, + IndexValueTuple<${INDEX_TYPE}, ${DATA_TYPE}>* _rs, + const typename SB_Handle::event_t& dependencies); -template typename SB_Handle::event_t _iamax(SB_Handle &sb_handle, - ${INDEX_TYPE} _N, - const ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, - IndexValueTuple<${INDEX_TYPE}, ${container_t1}>* _rs, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _iamax( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, const ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, + IndexValueTuple<${INDEX_TYPE}, ${DATA_TYPE}>* _rs, + const typename SB_Handle::event_t& dependencies); #endif } // namespace internal diff --git a/src/interface/blas1/iamax_return.cpp.in b/src/interface/blas1/iamax_return.cpp.in index 85f5c3714..f3eab8477 100644 --- a/src/interface/blas1/iamax_return.cpp.in +++ b/src/interface/blas1/iamax_return.cpp.in @@ -39,17 +39,19 @@ namespace internal { * @param _incx Increment in X axis * @param SB_Handle sb_handle */ -template ${INDEX_TYPE} _iamax(SB_Handle &sb_handle, ${INDEX_TYPE} _N, - BufferIterator<${container_t0}> _vx, ${INCREMENT_TYPE} _incx, +template ${INDEX_TYPE} _iamax(SB_Handle& sb_handle, ${INDEX_TYPE} _N, + BufferIterator<${DATA_TYPE}> _vx, + ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM -template ${INDEX_TYPE} _iamax(SB_Handle &sb_handle, ${INDEX_TYPE} _N, - ${container_t0}* _vx, ${INCREMENT_TYPE} _incx, +template ${INDEX_TYPE} _iamax(SB_Handle& sb_handle, ${INDEX_TYPE} _N, + ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& dependencies); -template ${INDEX_TYPE} _iamax(SB_Handle &sb_handle, ${INDEX_TYPE} _N, - const ${container_t0}* _vx, ${INCREMENT_TYPE} _incx, +template ${INDEX_TYPE} _iamax(SB_Handle& sb_handle, ${INDEX_TYPE} _N, + const ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& dependencies); #endif diff --git a/src/interface/blas1/iamin.cpp.in b/src/interface/blas1/iamin.cpp.in index 82a0e4fda..6fcd01798 100644 --- a/src/interface/blas1/iamin.cpp.in +++ b/src/interface/blas1/iamin.cpp.in @@ -39,27 +39,24 @@ namespace internal { * @param _vx VectorView * @param _incx Increment in X axis */ -template typename SB_Handle::event_t _iamin(SB_Handle &sb_handle, - ${INDEX_TYPE} _N, - BufferIterator<${container_t0}> _vx, - ${INCREMENT_TYPE} _incx, - BufferIterator> _rs, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _iamin( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, BufferIterator<${DATA_TYPE}> _vx, + ${INCREMENT_TYPE} _incx, + BufferIterator> _rs, + const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM -template typename SB_Handle::event_t _iamin(SB_Handle &sb_handle, - ${INDEX_TYPE} _N, - ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, - IndexValueTuple<${INDEX_TYPE}, ${container_t1}>* _rs, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _iamin( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, + IndexValueTuple<${INDEX_TYPE}, ${DATA_TYPE}>* _rs, + const typename SB_Handle::event_t& dependencies); -template typename SB_Handle::event_t _iamin(SB_Handle &sb_handle, - ${INDEX_TYPE} _N, - const ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, - IndexValueTuple<${INDEX_TYPE}, ${container_t1}>* _rs, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _iamin( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, const ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, + IndexValueTuple<${INDEX_TYPE}, ${DATA_TYPE}>* _rs, + const typename SB_Handle::event_t& dependencies); #endif diff --git a/src/interface/blas1/iamin_return.cpp.in b/src/interface/blas1/iamin_return.cpp.in index 7299b3afd..f3d39498f 100644 --- a/src/interface/blas1/iamin_return.cpp.in +++ b/src/interface/blas1/iamin_return.cpp.in @@ -39,17 +39,19 @@ namespace internal { * @param _vx VectorView * @param _incx Increment in X axis */ -template ${INDEX_TYPE} _iamin(SB_Handle &sb_handle, ${INDEX_TYPE} _N, - BufferIterator<${container_t0}> _vx, ${INCREMENT_TYPE} _incx, +template ${INDEX_TYPE} _iamin(SB_Handle& sb_handle, ${INDEX_TYPE} _N, + BufferIterator<${DATA_TYPE}> _vx, + ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM -template ${INDEX_TYPE} _iamin(SB_Handle &sb_handle, ${INDEX_TYPE} _N, - ${container_t0}* _vx, ${INCREMENT_TYPE} _incx, +template ${INDEX_TYPE} _iamin(SB_Handle& sb_handle, ${INDEX_TYPE} _N, + ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& dependencies); -template ${INDEX_TYPE} _iamin(SB_Handle &sb_handle, ${INDEX_TYPE} _N, - const ${container_t0}* _vx, ${INCREMENT_TYPE} _incx, +template ${INDEX_TYPE} _iamin(SB_Handle& sb_handle, ${INDEX_TYPE} _N, + const ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& dependencies); #endif diff --git a/src/interface/blas1/nrm2.cpp.in b/src/interface/blas1/nrm2.cpp.in index 478ec8541..88ac60595 100644 --- a/src/interface/blas1/nrm2.cpp.in +++ b/src/interface/blas1/nrm2.cpp.in @@ -39,27 +39,21 @@ namespace internal { * @param _vx VectorView * @param _incx Increment in X axis */ -template typename SB_Handle::event_t _nrm2(SB_Handle &sb_handle, - ${INDEX_TYPE} _N, - BufferIterator<${container_t0}> _vx, - ${INCREMENT_TYPE} _incx, - BufferIterator<${container_t1}> _rs, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _nrm2( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, BufferIterator<${DATA_TYPE}> _vx, + ${INCREMENT_TYPE} _incx, BufferIterator<${DATA_TYPE}> _rs, + const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM -template typename SB_Handle::event_t _nrm2(SB_Handle &sb_handle, - ${INDEX_TYPE} _N, - ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, - ${container_t1}* _rs, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _nrm2( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _rs, + const typename SB_Handle::event_t& dependencies); -template typename SB_Handle::event_t _nrm2(SB_Handle &sb_handle, - ${INDEX_TYPE} _N, - const ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, - ${container_t1}* _rs, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _nrm2( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, const ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _rs, + const typename SB_Handle::event_t& dependencies); #endif } // namespace internal diff --git a/src/interface/blas1/nrm2_return.cpp.in b/src/interface/blas1/nrm2_return.cpp.in index e176eca11..9a8ca35c7 100644 --- a/src/interface/blas1/nrm2_return.cpp.in +++ b/src/interface/blas1/nrm2_return.cpp.in @@ -40,17 +40,17 @@ namespace internal { * @param _vx VectorView * @param _incx Increment in X axis */ -template typename ValueType>::type _nrm2( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, BufferIterator<${container_t0}> _vx, +template typename ValueType>::type _nrm2( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, BufferIterator<${DATA_TYPE}> _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM -template typename ValueType<${container_t0}*>::type _nrm2( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, ${container_t0}* _vx, +template typename ValueType<${DATA_TYPE}*>::type _nrm2( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& dependencies); -template typename ValueType<${container_t0}*>::type _nrm2( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, const ${container_t0}* _vx, +template typename ValueType<${DATA_TYPE}*>::type _nrm2( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, const ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& dependencies); #endif diff --git a/src/interface/blas1/rot.cpp.in b/src/interface/blas1/rot.cpp.in index 4eed6a041..885106c3f 100644 --- a/src/interface/blas1/rot.cpp.in +++ b/src/interface/blas1/rot.cpp.in @@ -44,22 +44,18 @@ namespace internal { * @param _cos cosine * @param _N data size */ -template typename SB_Handle::event_t _rot(SB_Handle &sb_handle, - ${INDEX_TYPE} _N, BufferIterator<${container_t0}> _vx, - ${INCREMENT_TYPE} _incx, - BufferIterator<${container_t1}> _vy, - ${INCREMENT_TYPE} _incy, - ${DATA_TYPE} _cos, ${DATA_TYPE} _sin, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _rot( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, BufferIterator<${DATA_TYPE}> _vx, + ${INCREMENT_TYPE} _incx, BufferIterator<${DATA_TYPE}> _vy, + ${INCREMENT_TYPE} _incy, ${DATA_TYPE} _cos, ${DATA_TYPE} _sin, + const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM -template typename SB_Handle::event_t _rot(SB_Handle &sb_handle, - ${INDEX_TYPE} _N, ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, - ${container_t1}* _vy, - ${INCREMENT_TYPE} _incy, - ${DATA_TYPE} _cos, ${DATA_TYPE} _sin, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _rot( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, + ${DATA_TYPE} _cos, ${DATA_TYPE} _sin, + const typename SB_Handle::event_t& dependencies); #endif } // namespace internal diff --git a/src/interface/blas1/rotg.cpp.in b/src/interface/blas1/rotg.cpp.in index bd8bd3e67..f5a25ad97 100644 --- a/src/interface/blas1/rotg.cpp.in +++ b/src/interface/blas1/rotg.cpp.in @@ -50,20 +50,17 @@ namespace internal { * @param s[out] Buffer holding the parameter s. * @return Vector of events to wait for. */ -template typename SB_Handle::event_t _rotg(SB_Handle &sb_handle, - BufferIterator<${container_t0}> a, - BufferIterator<${container_t1}> b, - BufferIterator<${container_t2}> c, - BufferIterator<${container_t3}> s, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _rotg( + SB_Handle& sb_handle, BufferIterator<${DATA_TYPE}> a, + BufferIterator<${DATA_TYPE}> b, BufferIterator<${DATA_TYPE}> c, + BufferIterator<${DATA_TYPE}> s, + const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM -template typename SB_Handle::event_t _rotg(SB_Handle &sb_handle, - ${container_t0}* a, - ${container_t1}* b, - ${container_t2}* c, - ${container_t3}* s, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _rotg( + SB_Handle& sb_handle, ${DATA_TYPE} * a, ${DATA_TYPE} * b, + ${DATA_TYPE} * c, ${DATA_TYPE} * s, + const typename SB_Handle::event_t& dependencies); #endif } // namespace internal diff --git a/src/interface/blas1/rotg_return.cpp.in b/src/interface/blas1/rotg_return.cpp.in index 4daf106d0..d6becb83f 100644 --- a/src/interface/blas1/rotg_return.cpp.in +++ b/src/interface/blas1/rotg_return.cpp.in @@ -45,7 +45,8 @@ namespace internal { * @param c[out] Scalar representing the output c. * @param s[out] Scalar representing the output s. */ -template void _rotg(SB_Handle &sb_handle, ${DATA_TYPE} & a, ${DATA_TYPE} & b, - ${DATA_TYPE} & c, ${DATA_TYPE} & s, const typename SB_Handle::event_t& dependencies); +template void _rotg(SB_Handle& sb_handle, ${DATA_TYPE} & a, ${DATA_TYPE} & b, + ${DATA_TYPE} & c, ${DATA_TYPE} & s, + const typename SB_Handle::event_t& dependencies); } // namespace internal } // namespace blas diff --git a/src/interface/blas1/rotm.cpp.in b/src/interface/blas1/rotm.cpp.in index 81cba803f..455ad31d6 100644 --- a/src/interface/blas1/rotm.cpp.in +++ b/src/interface/blas1/rotm.cpp.in @@ -66,20 +66,22 @@ namespace internal { * @return Vector of events to wait for. */ template typename SB_Handle::event_t _rotm( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, BufferIterator<${container_t0}> _vx, - ${INCREMENT_TYPE} _incx, BufferIterator<${container_t1}> _vy, ${INCREMENT_TYPE} _incy, - BufferIterator<${container_t2}> _param, const typename SB_Handle::event_t& dependencies); + SB_Handle& sb_handle, ${INDEX_TYPE} _N, BufferIterator<${DATA_TYPE}> _vx, + ${INCREMENT_TYPE} _incx, BufferIterator<${DATA_TYPE}> _vy, + ${INCREMENT_TYPE} _incy, BufferIterator<${DATA_TYPE}> _param, + const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _rotm( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, ${container_t1}* _vy, ${INCREMENT_TYPE} _incy, - ${container_t2}* _param, const typename SB_Handle::event_t& dependencies); + SB_Handle& sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, + ${DATA_TYPE} * _param, const typename SB_Handle::event_t& dependencies); template typename SB_Handle::event_t _rotm( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, ${container_t1}* _vy, ${INCREMENT_TYPE} _incy, - const ${container_t2}* _param, const typename SB_Handle::event_t& dependencies); + SB_Handle& sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, + const ${DATA_TYPE} * _param, + const typename SB_Handle::event_t& dependencies); #endif } // namespace internal diff --git a/src/interface/blas1/rotmg.cpp.in b/src/interface/blas1/rotmg.cpp.in index 29d385238..573985e86 100644 --- a/src/interface/blas1/rotmg.cpp.in +++ b/src/interface/blas1/rotmg.cpp.in @@ -67,14 +67,15 @@ namespace internal { * @return Vector of events to wait for. */ template typename SB_Handle::event_t _rotmg( - SB_Handle &sb_handle, BufferIterator<${container_t0}> _d1, BufferIterator<${container_t1}> _d2, - BufferIterator<${container_t2}> _x1, BufferIterator<${container_t3}> _y1, BufferIterator<${container_t4}> _param, + SB_Handle& sb_handle, BufferIterator<${DATA_TYPE}> _d1, + BufferIterator<${DATA_TYPE}> _d2, BufferIterator<${DATA_TYPE}> _x1, + BufferIterator<${DATA_TYPE}> _y1, BufferIterator<${DATA_TYPE}> _param, const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _rotmg( - SB_Handle &sb_handle, ${container_t0}* _d1, ${container_t1}* _d2, - ${container_t2}* _x1, ${container_t3}* _y1, ${container_t4}* _param, + SB_Handle& sb_handle, ${DATA_TYPE} * _d1, ${DATA_TYPE} * _d2, + ${DATA_TYPE} * _x1, ${DATA_TYPE} * _y1, ${DATA_TYPE} * _param, const typename SB_Handle::event_t& dependencies); #endif diff --git a/src/interface/blas1/scal.cpp.in b/src/interface/blas1/scal.cpp.in index 951833f7c..9c2578c18 100644 --- a/src/interface/blas1/scal.cpp.in +++ b/src/interface/blas1/scal.cpp.in @@ -39,20 +39,16 @@ namespace internal { * @param _vx VectorView * @param _incx Increment in X axis */ -template typename SB_Handle::event_t _scal(SB_Handle &sb_handle, - ${INDEX_TYPE} _N, - ${DATA_TYPE} _alpha, - BufferIterator<${container_t0}> _vx, - ${INCREMENT_TYPE} _incx, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _scal( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, + BufferIterator<${DATA_TYPE}> _vx, ${INCREMENT_TYPE} _incx, + const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM -template typename SB_Handle::event_t _scal(SB_Handle &sb_handle, - ${INDEX_TYPE} _N, - ${DATA_TYPE} _alpha, - ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _scal( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, + ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, + const typename SB_Handle::event_t& dependencies); #endif } // namespace internal diff --git a/src/interface/blas1/sdsdot.cpp.in b/src/interface/blas1/sdsdot.cpp.in index 2aca7cf18..4c6e1cb87 100644 --- a/src/interface/blas1/sdsdot.cpp.in +++ b/src/interface/blas1/sdsdot.cpp.in @@ -54,15 +54,17 @@ namespace internal { * @return Vector of events to wait for. */ template typename SB_Handle::event_t _sdsdot( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, float sb, BufferIterator<${container_t0}> _vx, - ${INCREMENT_TYPE} _incx, BufferIterator<${container_t1}> _vy, ${INCREMENT_TYPE} _incy, - BufferIterator<${container_t2}> _rs, const typename SB_Handle::event_t& dependencies); + SB_Handle& sb_handle, ${INDEX_TYPE} _N, float sb, + BufferIterator<${DATA_TYPE}> _vx, ${INCREMENT_TYPE} _incx, + BufferIterator<${DATA_TYPE}> _vy, ${INCREMENT_TYPE} _incy, + BufferIterator<${DATA_TYPE}> _rs, + const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _sdsdot( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, float sb, ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, ${container_t1}* _vy, ${INCREMENT_TYPE} _incy, - ${container_t2}* _rs, const typename SB_Handle::event_t& dependencies); + SB_Handle& sb_handle, ${INDEX_TYPE} _N, float sb, ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, + ${DATA_TYPE} * _rs, const typename SB_Handle::event_t& dependencies); #endif } // namespace internal diff --git a/src/interface/blas1/sdsdot_return.cpp.in b/src/interface/blas1/sdsdot_return.cpp.in index b7e57f1c3..28004d042 100644 --- a/src/interface/blas1/sdsdot_return.cpp.in +++ b/src/interface/blas1/sdsdot_return.cpp.in @@ -53,15 +53,16 @@ namespace internal { * @param _rs Output buffer * @return Vector of events to wait for. */ -template typename ValueType>::type _sdsdot( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, float sb, BufferIterator<${container_t0}> _vx, - ${INCREMENT_TYPE} _incx, BufferIterator<${container_t1}> _vy, ${INCREMENT_TYPE} _incy, +template typename ValueType>::type _sdsdot( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, float sb, + BufferIterator<${DATA_TYPE}> _vx, ${INCREMENT_TYPE} _incx, + BufferIterator<${DATA_TYPE}> _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM -template typename ValueType<${container_t0}>::type _sdsdot( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, float sb, ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, ${container_t1}* _vy, ${INCREMENT_TYPE} _incy, +template typename ValueType<${DATA_TYPE}>::type _sdsdot( + SB_Handle& sb_handle, ${INDEX_TYPE} _N, float sb, ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& dependencies); #endif diff --git a/src/interface/blas1/swap.cpp.in b/src/interface/blas1/swap.cpp.in index 4de9508ee..34cac5d69 100644 --- a/src/interface/blas1/swap.cpp.in +++ b/src/interface/blas1/swap.cpp.in @@ -42,14 +42,14 @@ namespace internal { * @param _incy Increment in Y axis */ template typename SB_Handle::event_t _swap( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, BufferIterator<${container_t0}> _vx, - ${INCREMENT_TYPE} _incx, BufferIterator<${container_t1}> _vy, ${INCREMENT_TYPE} _incy, - const typename SB_Handle::event_t& dependencies); + SB_Handle& sb_handle, ${INDEX_TYPE} _N, BufferIterator<${DATA_TYPE}> _vx, + ${INCREMENT_TYPE} _incx, BufferIterator<${DATA_TYPE}> _vy, + ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _swap( - SB_Handle &sb_handle, ${INDEX_TYPE} _N, ${container_t0}* _vx, - ${INCREMENT_TYPE} _incx, ${container_t1}* _vy, ${INCREMENT_TYPE} _incy, + SB_Handle& sb_handle, ${INDEX_TYPE} _N, ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& dependencies); #endif diff --git a/src/interface/blas2/CMakeLists.txt b/src/interface/blas2/CMakeLists.txt index 01c2e9513..4f99dc45a 100644 --- a/src/interface/blas2/CMakeLists.txt +++ b/src/interface/blas2/CMakeLists.txt @@ -24,18 +24,18 @@ # * # **************************************************************************/ #blas2 -generate_blas_ternary_objects(blas2 gbmv) -generate_blas_ternary_objects(blas2 gemv) -generate_blas_ternary_objects(blas2 ger) -generate_blas_ternary_objects(blas2 sbmv) -generate_blas_ternary_objects(blas2 spmv) -generate_blas_ternary_objects(blas2 symv) -generate_blas_ternary_objects(blas2 syr2) -generate_blas_binary_objects(blas2 syr) -generate_blas_binary_objects(blas2 tbmv) -generate_blas_binary_objects(blas2 tpmv) -generate_blas_binary_objects(blas2 spr) -generate_blas_ternary_objects(blas2 spr2) -generate_blas_binary_objects(blas2 tbsv) -generate_blas_binary_objects(blas2 trmv) -generate_blas_binary_objects(blas2 trsv) +generate_blas_unary_objects(blas2 gbmv) +generate_blas_unary_objects(blas2 gemv) +generate_blas_unary_objects(blas2 ger) +generate_blas_unary_objects(blas2 sbmv) +generate_blas_unary_objects(blas2 spmv) +generate_blas_unary_objects(blas2 symv) +generate_blas_unary_objects(blas2 syr2) +generate_blas_unary_objects(blas2 syr) +generate_blas_unary_objects(blas2 tbmv) +generate_blas_unary_objects(blas2 tpmv) +generate_blas_unary_objects(blas2 spr) +generate_blas_unary_objects(blas2 spr2) +generate_blas_unary_objects(blas2 tbsv) +generate_blas_unary_objects(blas2 trmv) +generate_blas_unary_objects(blas2 trsv) diff --git a/src/interface/blas2/gbmv.cpp.in b/src/interface/blas2/gbmv.cpp.in index e2d11b74d..bc6c62d8a 100644 --- a/src/interface/blas2/gbmv.cpp.in +++ b/src/interface/blas2/gbmv.cpp.in @@ -32,24 +32,26 @@ namespace internal { template typename SB_Handle::event_t _gbmv( SB_Handle& sb_handle, char _trans, ${INDEX_TYPE} _M, ${INDEX_TYPE} _N, ${INDEX_TYPE} _KL, ${INDEX_TYPE} _KU, ${DATA_TYPE} _alpha, - BufferIterator<${container_t0}> _mA, ${INDEX_TYPE} _lda, BufferIterator<${container_t1}> _vx, - ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, BufferIterator<${container_t2}> _vy, + BufferIterator<${DATA_TYPE}> _mA, ${INDEX_TYPE} _lda, + BufferIterator<${DATA_TYPE}> _vx, ${INCREMENT_TYPE} _incx, + ${DATA_TYPE} _beta, BufferIterator<${DATA_TYPE}> _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& _dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _gbmv( SB_Handle& sb_handle, char _trans, ${INDEX_TYPE} _M, ${INDEX_TYPE} _N, ${INDEX_TYPE} _KL, ${INDEX_TYPE} _KU, ${DATA_TYPE} _alpha, - ${container_t0}* _mA, ${INDEX_TYPE} _lda, ${container_t1}* _vx, - ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, ${container_t2}* _vy, + ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& _dependencies); template typename SB_Handle::event_t _gbmv( SB_Handle& sb_handle, char _trans, ${INDEX_TYPE} _M, ${INDEX_TYPE} _N, ${INDEX_TYPE} _KL, ${INDEX_TYPE} _KU, ${DATA_TYPE} _alpha, - const ${container_t0}* _mA, ${INDEX_TYPE} _lda, const ${container_t1}* _vx, - ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, ${container_t2}* _vy, - ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& _dependencies); + const ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, + const ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, + ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, + const typename SB_Handle::event_t& _dependencies); #endif } // namespace internal diff --git a/src/interface/blas2/gemv.cpp.in b/src/interface/blas2/gemv.cpp.in index 7471cb861..2c6b980dd 100644 --- a/src/interface/blas2/gemv.cpp.in +++ b/src/interface/blas2/gemv.cpp.in @@ -63,33 +63,35 @@ namespace internal { */ template typename SB_Handle::event_t _gemv( SB_Handle& sb_handle, char _trans, ${INDEX_TYPE} _M, ${INDEX_TYPE} _N, - ${DATA_TYPE} _alpha, BufferIterator<${container_t0}> _mA, ${INDEX_TYPE} _lda, - BufferIterator<${container_t1}> _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, - BufferIterator<${container_t2}> _vy, ${INCREMENT_TYPE} _incy, + ${DATA_TYPE} _alpha, BufferIterator<${DATA_TYPE}> _mA, + ${INDEX_TYPE} _lda, BufferIterator<${DATA_TYPE}> _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, + BufferIterator<${DATA_TYPE}> _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& _dependencies); #ifdef BLAS_ENABLE_CONST_INPUT template typename SB_Handle::event_t _gemv( SB_Handle& sb_handle, char _trans, ${INDEX_TYPE} _M, ${INDEX_TYPE} _N, - ${DATA_TYPE} _alpha, BufferIterator<${container_t0} const> _mA, ${INDEX_TYPE} _lda, - BufferIterator<${container_t1} const> _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, - BufferIterator<${container_t2}> _vy, ${INCREMENT_TYPE} _incy, + ${DATA_TYPE} _alpha, BufferIterator<${DATA_TYPE} const> _mA, + ${INDEX_TYPE} _lda, BufferIterator<${DATA_TYPE} const> _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, + BufferIterator<${DATA_TYPE}> _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& _dependencies); #endif #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _gemv( SB_Handle& sb_handle, char _trans, ${INDEX_TYPE} _M, ${INDEX_TYPE} _N, - ${DATA_TYPE} _alpha, ${container_t0}* _mA, ${INDEX_TYPE} _lda, - ${container_t1}* _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, - ${container_t2}* _vy, ${INCREMENT_TYPE} _incy, + ${DATA_TYPE} _alpha, ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, + ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, + ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& _dependencies); template typename SB_Handle::event_t _gemv( SB_Handle& sb_handle, char _trans, ${INDEX_TYPE} _M, ${INDEX_TYPE} _N, - ${DATA_TYPE} _alpha, const ${container_t0}* _mA, ${INDEX_TYPE} _lda, - const ${container_t1}* _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, - ${container_t2}* _vy, ${INCREMENT_TYPE} _incy, + ${DATA_TYPE} _alpha, const ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, + const ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, + ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& _dependencies); #endif diff --git a/src/interface/blas2/ger.cpp.in b/src/interface/blas2/ger.cpp.in index 8a2979814..8a3f11bcf 100644 --- a/src/interface/blas2/ger.cpp.in +++ b/src/interface/blas2/ger.cpp.in @@ -32,21 +32,22 @@ namespace internal { template typename SB_Handle::event_t _ger( SB_Handle& sb_handle, ${INDEX_TYPE} _M, ${INDEX_TYPE} _N, - ${DATA_TYPE} _alpha, BufferIterator<${container_t0}> _vx, ${INCREMENT_TYPE} _incx, - BufferIterator<${container_t1}> _vy, ${INCREMENT_TYPE} _incy, BufferIterator<${container_t2}> _mA, + ${DATA_TYPE} _alpha, BufferIterator<${DATA_TYPE}> _vx, + ${INCREMENT_TYPE} _incx, BufferIterator<${DATA_TYPE}> _vy, + ${INCREMENT_TYPE} _incy, BufferIterator<${DATA_TYPE}> _mA, ${INDEX_TYPE} _lda, const typename SB_Handle::event_t& _dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _ger( SB_Handle& sb_handle, ${INDEX_TYPE} _M, ${INDEX_TYPE} _N, - ${DATA_TYPE} _alpha, ${container_t0}* _vx, ${INCREMENT_TYPE} _incx, - ${container_t1}* _vy, ${INCREMENT_TYPE} _incy, ${container_t2}* _mA, + ${DATA_TYPE} _alpha, ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, + ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, const typename SB_Handle::event_t& _dependencies); template typename SB_Handle::event_t _ger( SB_Handle& sb_handle, ${INDEX_TYPE} _M, ${INDEX_TYPE} _N, - ${DATA_TYPE} _alpha, const ${container_t0}* _vx, ${INCREMENT_TYPE} _incx, - const ${container_t1}* _vy, ${INCREMENT_TYPE} _incy, ${container_t2}* _mA, + ${DATA_TYPE} _alpha, const ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, + const ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, const typename SB_Handle::event_t& _dependencies); #endif diff --git a/src/interface/blas2/sbmv.cpp.in b/src/interface/blas2/sbmv.cpp.in index 073bca2e7..cc3d42df4 100644 --- a/src/interface/blas2/sbmv.cpp.in +++ b/src/interface/blas2/sbmv.cpp.in @@ -31,24 +31,25 @@ namespace blas { namespace internal { template typename SB_Handle::event_t _sbmv( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, - ${DATA_TYPE} _alpha, BufferIterator<${container_t0}> _mA, ${INDEX_TYPE} _lda, - BufferIterator<${container_t1}> _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, - BufferIterator<${container_t2}> _vy, ${INCREMENT_TYPE} _incy, + ${DATA_TYPE} _alpha, BufferIterator<${DATA_TYPE}> _mA, + ${INDEX_TYPE} _lda, BufferIterator<${DATA_TYPE}> _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, + BufferIterator<${DATA_TYPE}> _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& _dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _sbmv( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, - ${DATA_TYPE} _alpha, ${container_t0}* _mA, ${INDEX_TYPE} _lda, - ${container_t1}* _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, - ${container_t2}* _vy, ${INCREMENT_TYPE} _incy, + ${DATA_TYPE} _alpha, ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, + ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, + ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& _dependencies); template typename SB_Handle::event_t _sbmv( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, - ${DATA_TYPE} _alpha, const ${container_t0}* _mA, ${INDEX_TYPE} _lda, - const ${container_t1}* _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, - ${container_t2}* _vy, ${INCREMENT_TYPE} _incy, + ${DATA_TYPE} _alpha, const ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, + const ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, + ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& _dependencies); #endif diff --git a/src/interface/blas2/spmv.cpp.in b/src/interface/blas2/spmv.cpp.in index ff8be406b..9c53cc9b2 100644 --- a/src/interface/blas2/spmv.cpp.in +++ b/src/interface/blas2/spmv.cpp.in @@ -30,22 +30,23 @@ namespace blas { namespace internal { template typename SB_Handle::event_t _spmv( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - BufferIterator<${container_t0}> _mA, BufferIterator<${container_t1}> _vx, ${INCREMENT_TYPE} _incx, - ${DATA_TYPE} _beta, BufferIterator<${container_t2}> _vy, ${INCREMENT_TYPE} _incy, + BufferIterator<${DATA_TYPE}> _mA, BufferIterator<${DATA_TYPE}> _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, + BufferIterator<${DATA_TYPE}> _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& _dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _spmv( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - ${container_t0}* _mA, ${container_t1}* _vx, ${INCREMENT_TYPE} _incx, - ${DATA_TYPE} _beta, ${container_t2}* _vy, ${INCREMENT_TYPE} _incy, + ${DATA_TYPE} * _mA, ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, + ${DATA_TYPE} _beta, ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& _dependencies); template typename SB_Handle::event_t _spmv( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - const ${container_t0}* _mA, const ${container_t1}* _vx, ${INCREMENT_TYPE} _incx, - ${DATA_TYPE} _beta, ${container_t2}* _vy, ${INCREMENT_TYPE} _incy, - const typename SB_Handle::event_t& _dependencies); + const ${DATA_TYPE} * _mA, const ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, ${DATA_TYPE} * _vy, + ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& _dependencies); #endif } // namespace internal diff --git a/src/interface/blas2/spr.cpp.in b/src/interface/blas2/spr.cpp.in index eac14083e..12e93a5dc 100644 --- a/src/interface/blas2/spr.cpp.in +++ b/src/interface/blas2/spr.cpp.in @@ -32,19 +32,20 @@ namespace internal { template typename SB_Handle::event_t _spr( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - BufferIterator<${container_t0}> _vx, ${INCREMENT_TYPE} _incx, BufferIterator<${container_t1}> _mPA, + BufferIterator<${DATA_TYPE}> _vx, ${INCREMENT_TYPE} _incx, + BufferIterator<${DATA_TYPE}> _mPA, const typename SB_Handle::event_t& _dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _spr( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - ${container_t0}* _vx, ${INCREMENT_TYPE} _incx, ${container_t1}* _mPA, + ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _mPA, const typename SB_Handle::event_t& _dependencies); template typename SB_Handle::event_t _spr( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - const ${container_t0}* _vx, ${INCREMENT_TYPE} _incx, ${container_t1}* _mPA, - const typename SB_Handle::event_t& _dependencies); + const ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, + ${DATA_TYPE} * _mPA, const typename SB_Handle::event_t& _dependencies); #endif } // namespace internal diff --git a/src/interface/blas2/spr2.cpp.in b/src/interface/blas2/spr2.cpp.in index a0959814b..3c5f1250b 100644 --- a/src/interface/blas2/spr2.cpp.in +++ b/src/interface/blas2/spr2.cpp.in @@ -31,22 +31,23 @@ namespace internal { template typename SB_Handle::event_t _spr2( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - BufferIterator<${container_t0}> _vx, ${INCREMENT_TYPE} _incx, BufferIterator<${container_t1}> _vy, - ${INCREMENT_TYPE} _incy, BufferIterator<${container_t2}> _mPA, + BufferIterator<${DATA_TYPE}> _vx, ${INCREMENT_TYPE} _incx, + BufferIterator<${DATA_TYPE}> _vy, ${INCREMENT_TYPE} _incy, + BufferIterator<${DATA_TYPE}> _mPA, const typename SB_Handle::event_t& _dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _spr2( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - ${container_t0}* _vx, ${INCREMENT_TYPE} _incx, ${container_t1}* _vy, - ${INCREMENT_TYPE} _incy, ${container_t2}* _mPA, + ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _vy, + ${INCREMENT_TYPE} _incy, ${DATA_TYPE} * _mPA, const typename SB_Handle::event_t& _dependencies); template typename SB_Handle::event_t _spr2( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - const ${container_t0}* _vx, ${INCREMENT_TYPE} _incx, const ${container_t1}* _vy, - ${INCREMENT_TYPE} _incy, ${container_t2}* _mPA, - const typename SB_Handle::event_t& _dependencies); + const ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, + const ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, + ${DATA_TYPE} * _mPA, const typename SB_Handle::event_t& _dependencies); #endif } // namespace internal diff --git a/src/interface/blas2/symv.cpp.in b/src/interface/blas2/symv.cpp.in index 504822abe..225411025 100644 --- a/src/interface/blas2/symv.cpp.in +++ b/src/interface/blas2/symv.cpp.in @@ -32,22 +32,24 @@ namespace internal { template typename SB_Handle::event_t _symv( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - BufferIterator<${container_t0}> _mA, ${INDEX_TYPE} _lda, BufferIterator<${container_t1}> _vx, - ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, BufferIterator<${container_t2}> _vy, + BufferIterator<${DATA_TYPE}> _mA, ${INDEX_TYPE} _lda, + BufferIterator<${DATA_TYPE}> _vx, ${INCREMENT_TYPE} _incx, + ${DATA_TYPE} _beta, BufferIterator<${DATA_TYPE}> _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& _dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _symv( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - ${container_t0}* _mA, ${INDEX_TYPE} _lda, ${container_t1}* _vx, - ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, ${container_t2}* _vy, + ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, ${DATA_TYPE} * _vx, + ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& _dependencies); template typename SB_Handle::event_t _symv( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - const ${container_t0}* _mA, ${INDEX_TYPE} _lda, const ${container_t1}* _vx, - ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, ${container_t2}* _vy, - ${INCREMENT_TYPE} _incy, const typename SB_Handle::event_t& _dependencies); + const ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, + const ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} _beta, + ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, + const typename SB_Handle::event_t& _dependencies); #endif } // namespace internal diff --git a/src/interface/blas2/syr.cpp.in b/src/interface/blas2/syr.cpp.in index 3e5b61b80..358dfbff6 100644 --- a/src/interface/blas2/syr.cpp.in +++ b/src/interface/blas2/syr.cpp.in @@ -32,18 +32,19 @@ namespace internal { template typename SB_Handle::event_t _syr( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - BufferIterator<${container_t0}> _vx, ${INCREMENT_TYPE} _incx, BufferIterator<${container_t1}> _mA, - ${INDEX_TYPE} _lda, const typename SB_Handle::event_t& _dependencies); + BufferIterator<${DATA_TYPE}> _vx, ${INCREMENT_TYPE} _incx, + BufferIterator<${DATA_TYPE}> _mA, ${INDEX_TYPE} _lda, + const typename SB_Handle::event_t& _dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _syr( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - ${container_t0}* _vx, ${INCREMENT_TYPE} _incx, ${container_t1}* _mA, + ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, const typename SB_Handle::event_t& _dependencies); template typename SB_Handle::event_t _syr( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - const ${container_t0}* _vx, ${INCREMENT_TYPE} _incx, ${container_t1}* _mA, + const ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, const typename SB_Handle::event_t& _dependencies); #endif diff --git a/src/interface/blas2/syr2.cpp.in b/src/interface/blas2/syr2.cpp.in index 68c459f34..070977b8e 100644 --- a/src/interface/blas2/syr2.cpp.in +++ b/src/interface/blas2/syr2.cpp.in @@ -32,23 +32,24 @@ namespace internal { template typename SB_Handle::event_t _syr2( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - BufferIterator<${container_t0}> _vx, ${INCREMENT_TYPE} _incx, BufferIterator<${container_t1}> _vy, - ${INCREMENT_TYPE} _incy, BufferIterator<${container_t2}> _mA, ${INDEX_TYPE} _lda, + BufferIterator<${DATA_TYPE}> _vx, ${INCREMENT_TYPE} _incx, + BufferIterator<${DATA_TYPE}> _vy, ${INCREMENT_TYPE} _incy, + BufferIterator<${DATA_TYPE}> _mA, ${INDEX_TYPE} _lda, const typename SB_Handle::event_t& _dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _syr2( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - ${container_t0}* _vx, ${INCREMENT_TYPE} _incx, ${container_t1}* _vy, - ${INCREMENT_TYPE} _incy, ${container_t2}* _mA, ${INDEX_TYPE} _lda, + ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, ${DATA_TYPE} * _vy, + ${INCREMENT_TYPE} _incy, ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, const typename SB_Handle::event_t& _dependencies); template typename SB_Handle::event_t _syr2( SB_Handle& sb_handle, char _Uplo, ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, - const ${container_t0}* _vx, ${INCREMENT_TYPE} _incx, const ${container_t1}* _vy, - ${INCREMENT_TYPE} _incy, ${container_t2}* _mA, ${INDEX_TYPE} _lda, - const typename SB_Handle::event_t& _dependencies); + const ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, + const ${DATA_TYPE} * _vy, ${INCREMENT_TYPE} _incy, ${DATA_TYPE} * _mA, + ${INDEX_TYPE} _lda, const typename SB_Handle::event_t& _dependencies); #endif -} +} // namespace internal } // end namespace blas diff --git a/src/interface/blas2/tbmv.cpp.in b/src/interface/blas2/tbmv.cpp.in index 11b59290b..9aeb0192b 100644 --- a/src/interface/blas2/tbmv.cpp.in +++ b/src/interface/blas2/tbmv.cpp.in @@ -31,21 +31,21 @@ namespace blas { namespace internal { template typename SB_Handle::event_t _tbmv( SB_Handle& sb_handle, char _Uplo, char _trans, char _Diag, ${INDEX_TYPE} _N, - ${INDEX_TYPE} _K, BufferIterator<${container_t0}> _mA, ${INDEX_TYPE} _lda, - BufferIterator<${container_t1}> _vx, ${INCREMENT_TYPE} _incx, + ${INDEX_TYPE} _K, BufferIterator<${DATA_TYPE}> _mA, ${INDEX_TYPE} _lda, + BufferIterator<${DATA_TYPE}> _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& _dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _tbmv( SB_Handle& sb_handle, char _Uplo, char _trans, char _Diag, ${INDEX_TYPE} _N, - ${INDEX_TYPE} _K, ${container_t0}* _mA, ${INDEX_TYPE} _lda, - ${container_t1}* _vx, ${INCREMENT_TYPE} _incx, + ${INDEX_TYPE} _K, ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, + ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& _dependencies); template typename SB_Handle::event_t _tbmv( SB_Handle& sb_handle, char _Uplo, char _trans, char _Diag, ${INDEX_TYPE} _N, - ${INDEX_TYPE} _K, const ${container_t0}* _mA, ${INDEX_TYPE} _lda, - ${container_t1}* _vx, ${INCREMENT_TYPE} _incx, + ${INDEX_TYPE} _K, const ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, + ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& _dependencies); #endif diff --git a/src/interface/blas2/tbsv.cpp.in b/src/interface/blas2/tbsv.cpp.in index 8f51b25f6..37ed365b3 100644 --- a/src/interface/blas2/tbsv.cpp.in +++ b/src/interface/blas2/tbsv.cpp.in @@ -31,21 +31,21 @@ namespace blas { namespace internal { template typename SB_Handle::event_t _tbsv( SB_Handle& sb_handle, char _Uplo, char _trans, char _Diag, ${INDEX_TYPE} _N, - ${INDEX_TYPE} _K, BufferIterator<${container_t0}> _mA, ${INDEX_TYPE} _lda, - BufferIterator<${container_t1}> _vx, ${INCREMENT_TYPE} _incx, + ${INDEX_TYPE} _K, BufferIterator<${DATA_TYPE}> _mA, ${INDEX_TYPE} _lda, + BufferIterator<${DATA_TYPE}> _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& _dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _tbsv( SB_Handle& sb_handle, char _Uplo, char _trans, char _Diag, ${INDEX_TYPE} _N, - ${INDEX_TYPE} _K, ${container_t0}* _mA, ${INDEX_TYPE} _lda, - ${container_t1}* _vx, ${INCREMENT_TYPE} _incx, + ${INDEX_TYPE} _K, ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, + ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& _dependencies); template typename SB_Handle::event_t _tbsv( SB_Handle& sb_handle, char _Uplo, char _trans, char _Diag, ${INDEX_TYPE} _N, - ${INDEX_TYPE} _K, const ${container_t0}* _mA, ${INDEX_TYPE} _lda, - ${container_t1}* _vx, ${INCREMENT_TYPE} _incx, + ${INDEX_TYPE} _K, const ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, + ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& _dependencies); #endif diff --git a/src/interface/blas2/tpmv.cpp.in b/src/interface/blas2/tpmv.cpp.in index 5d4044ed1..57a6b0e90 100644 --- a/src/interface/blas2/tpmv.cpp.in +++ b/src/interface/blas2/tpmv.cpp.in @@ -29,18 +29,18 @@ namespace blas { namespace internal { template typename SB_Handle::event_t _tpmv( SB_Handle& sb_handle, char _Uplo, char _trans, char _Diag, ${INDEX_TYPE} _N, - BufferIterator<${container_t0}> _mA, BufferIterator<${container_t1}> _vx, ${INCREMENT_TYPE} _incx, - const typename SB_Handle::event_t& _dependencies); + BufferIterator<${DATA_TYPE}> _mA, BufferIterator<${DATA_TYPE}> _vx, + ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& _dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _tpmv( SB_Handle& sb_handle, char _Uplo, char _trans, char _Diag, ${INDEX_TYPE} _N, - ${container_t0}* _mA, ${container_t1}* _vx, ${INCREMENT_TYPE} _incx, + ${DATA_TYPE} * _mA, ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& _dependencies); template typename SB_Handle::event_t _tpmv( SB_Handle& sb_handle, char _Uplo, char _trans, char _Diag, ${INDEX_TYPE} _N, - const ${container_t0}* _mA, ${container_t1}* _vx, ${INCREMENT_TYPE} _incx, + const ${DATA_TYPE} * _mA, ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& _dependencies); #endif diff --git a/src/interface/blas2/trmv.cpp.in b/src/interface/blas2/trmv.cpp.in index 6fb61ca70..a5aa48762 100644 --- a/src/interface/blas2/trmv.cpp.in +++ b/src/interface/blas2/trmv.cpp.in @@ -32,18 +32,19 @@ namespace internal { template typename SB_Handle::event_t _trmv( SB_Handle& sb_handle, char _Uplo, char _trans, char _Diag, ${INDEX_TYPE} _N, - BufferIterator<${container_t0}> _mA, ${INDEX_TYPE} _lda, BufferIterator<${container_t1}> _vx, - ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& _dependencies); + BufferIterator<${DATA_TYPE}> _mA, ${INDEX_TYPE} _lda, + BufferIterator<${DATA_TYPE}> _vx, ${INCREMENT_TYPE} _incx, + const typename SB_Handle::event_t& _dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _trmv( SB_Handle& sb_handle, char _Uplo, char _trans, char _Diag, ${INDEX_TYPE} _N, - ${container_t0}* _mA, ${INDEX_TYPE} _lda, ${container_t1}* _vx, + ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& _dependencies); template typename SB_Handle::event_t _trmv( SB_Handle& sb_handle, char _Uplo, char _trans, char _Diag, ${INDEX_TYPE} _N, - const ${container_t0}* _mA, ${INDEX_TYPE} _lda, ${container_t1}* _vx, + const ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& _dependencies); #endif diff --git a/src/interface/blas2/trsv.cpp.in b/src/interface/blas2/trsv.cpp.in index 9ea58e137..786c73580 100644 --- a/src/interface/blas2/trsv.cpp.in +++ b/src/interface/blas2/trsv.cpp.in @@ -32,18 +32,19 @@ namespace internal { template typename SB_Handle::event_t _trsv( SB_Handle& sb_handle, char _Uplo, char _trans, char _Diag, ${INDEX_TYPE} _N, - BufferIterator<${container_t0}> _mA, ${INDEX_TYPE} _lda, BufferIterator<${container_t1}> _vx, - ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& _dependencies); + BufferIterator<${DATA_TYPE}> _mA, ${INDEX_TYPE} _lda, + BufferIterator<${DATA_TYPE}> _vx, ${INCREMENT_TYPE} _incx, + const typename SB_Handle::event_t& _dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _trsv( SB_Handle& sb_handle, char _Uplo, char _trans, char _Diag, ${INDEX_TYPE} _N, - ${container_t0}* _mA, ${INDEX_TYPE} _lda, ${container_t1}* _vx, + ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& _dependencies); template typename SB_Handle::event_t _trsv( SB_Handle& sb_handle, char _Uplo, char _trans, char _Diag, ${INDEX_TYPE} _N, - const ${container_t0}* _mA, ${INDEX_TYPE} _lda, ${container_t1}* _vx, + const ${DATA_TYPE} * _mA, ${INDEX_TYPE} _lda, ${DATA_TYPE} * _vx, ${INCREMENT_TYPE} _incx, const typename SB_Handle::event_t& _dependencies); #endif diff --git a/src/interface/blas3/CMakeLists.txt b/src/interface/blas3/CMakeLists.txt index 446b65359..dfa8883dd 100644 --- a/src/interface/blas3/CMakeLists.txt +++ b/src/interface/blas3/CMakeLists.txt @@ -24,6 +24,6 @@ # **************************************************************************/ #blas3 generate_blas_gemm_objects(blas3 gemm_launcher) -generate_blas_ternary_objects(blas3 gemm) -generate_blas_ternary_objects(blas3 symm) -generate_blas_binary_objects(blas3 trsm) +generate_blas_unary_objects(blas3 gemm) +generate_blas_unary_objects(blas3 symm) +generate_blas_unary_objects(blas3 trsm) diff --git a/src/interface/blas3/gemm.cpp.in b/src/interface/blas3/gemm.cpp.in index 54995ef45..7ff401ef9 100644 --- a/src/interface/blas3/gemm.cpp.in +++ b/src/interface/blas3/gemm.cpp.in @@ -32,99 +32,108 @@ 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, BufferIterator<${container_t0}> a_, - ${INDEX_TYPE} _lda, BufferIterator<${container_t1}> b_, ${INDEX_TYPE} _ldb, - ${DATA_TYPE} _beta, BufferIterator<${container_t2}> _C, ${INDEX_TYPE} _ldc, + ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, ${DATA_TYPE} _alpha, + BufferIterator<${DATA_TYPE}> a_, ${INDEX_TYPE} _lda, + BufferIterator<${DATA_TYPE}> b_, ${INDEX_TYPE} _ldb, ${DATA_TYPE} _beta, + BufferIterator<${DATA_TYPE}> _C, ${INDEX_TYPE} _ldc, const typename SB_Handle::event_t& _dependencies); #ifdef SB_ENABLE_USM 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, - const typename SB_Handle::event_t& _dependencies); + ${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, const typename SB_Handle::event_t& _dependencies); 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, const ${container_t0}* a_, - ${INDEX_TYPE} _lda, const ${container_t1}* b_, ${INDEX_TYPE} _ldb, - ${DATA_TYPE} _beta, ${container_t2}* _C, ${INDEX_TYPE} _ldc, - const typename SB_Handle::event_t& _dependencies); + ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, ${DATA_TYPE} _alpha, + const ${DATA_TYPE} * a_, ${INDEX_TYPE} _lda, const ${DATA_TYPE} * b_, + ${INDEX_TYPE} _ldb, ${DATA_TYPE} _beta, ${DATA_TYPE} * _C, + ${INDEX_TYPE} _ldc, const typename SB_Handle::event_t& _dependencies); #endif // 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, BufferIterator<${container_t0}> a_, - ${INDEX_TYPE} _lda, BufferIterator<${container_t1}> b_, ${INDEX_TYPE} _ldb, - ${DATA_TYPE} _beta, BufferIterator<${container_t2}> _C, ${INDEX_TYPE} _ldc, + ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, ${DATA_TYPE} _alpha, + BufferIterator<${DATA_TYPE}> a_, ${INDEX_TYPE} _lda, + BufferIterator<${DATA_TYPE}> b_, ${INDEX_TYPE} _ldb, ${DATA_TYPE} _beta, + BufferIterator<${DATA_TYPE}> _C, ${INDEX_TYPE} _ldc, ${INDEX_TYPE} batch_size, gemm_batch_type_t batch_type, const typename SB_Handle::event_t& _dependencies); #ifdef SB_ENABLE_USM 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, ${container_t1}* b_, ${INDEX_TYPE} _ldb, - ${DATA_TYPE} _beta, ${container_t2}* _C, ${INDEX_TYPE} _ldc, - ${INDEX_TYPE} batch_size, gemm_batch_type_t batch_type, + ${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, ${INDEX_TYPE} batch_size, gemm_batch_type_t batch_type, const typename SB_Handle::event_t& _dependencies); 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, const ${container_t0}* a_, - ${INDEX_TYPE} _lda, const ${container_t1}* b_, ${INDEX_TYPE} _ldb, - ${DATA_TYPE} _beta, ${container_t2}* _C, ${INDEX_TYPE} _ldc, - ${INDEX_TYPE} batch_size, gemm_batch_type_t batch_type, + ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, ${DATA_TYPE} _alpha, + const ${DATA_TYPE} * a_, ${INDEX_TYPE} _lda, const ${DATA_TYPE} * b_, + ${INDEX_TYPE} _ldb, ${DATA_TYPE} _beta, ${DATA_TYPE} * _C, + ${INDEX_TYPE} _ldc, ${INDEX_TYPE} batch_size, gemm_batch_type_t batch_type, const typename SB_Handle::event_t& _dependencies); #endif // strided batched gemm template typename SB_Handle::event_t _gemm_strided_batched( SB_Handle& sb_handle, char _TransA, char _TransB, ${INDEX_TYPE} _M, - ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, ${DATA_TYPE} _alpha, BufferIterator<${container_t0}> a_, - ${INDEX_TYPE} _lda, ${INDEX_TYPE} _stridea, BufferIterator<${container_t1}> b_, + ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, ${DATA_TYPE} _alpha, + BufferIterator<${DATA_TYPE}> a_, ${INDEX_TYPE} _lda, + ${INDEX_TYPE} _stridea, BufferIterator<${DATA_TYPE}> b_, ${INDEX_TYPE} _ldb, ${INDEX_TYPE} _strideb, ${DATA_TYPE} _beta, - BufferIterator<${container_t2}> _C, ${INDEX_TYPE} _ldc, ${INDEX_TYPE} _stridec, - ${INDEX_TYPE} batch_size, const typename SB_Handle::event_t& _dependencies); + BufferIterator<${DATA_TYPE}> _C, ${INDEX_TYPE} _ldc, + ${INDEX_TYPE} _stridec, ${INDEX_TYPE} batch_size, + const typename SB_Handle::event_t& _dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _gemm_strided_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} _stridec, - ${INDEX_TYPE} batch_size, const typename SB_Handle::event_t& _dependencies); + ${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, + const typename SB_Handle::event_t& _dependencies); template typename SB_Handle::event_t _gemm_strided_batched( SB_Handle& sb_handle, char _TransA, char _TransB, ${INDEX_TYPE} _M, - ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, ${DATA_TYPE} _alpha, const ${container_t0}* a_, - ${INDEX_TYPE} _lda, ${INDEX_TYPE} _stridea, const ${container_t1}* b_, - ${INDEX_TYPE} _ldb, ${INDEX_TYPE} _strideb, ${DATA_TYPE} _beta, - ${container_t2}* _C, ${INDEX_TYPE} _ldc, ${INDEX_TYPE} _stridec, - ${INDEX_TYPE} batch_size, const typename SB_Handle::event_t& _dependencies); + ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, ${DATA_TYPE} _alpha, + const ${DATA_TYPE} * a_, ${INDEX_TYPE} _lda, ${INDEX_TYPE} _stridea, + const ${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, + const typename SB_Handle::event_t& _dependencies); #endif - #ifdef BLAS_ENABLE_CONST_INPUT 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, BufferIterator<${container_t0} const> a_, - ${INDEX_TYPE} _lda, BufferIterator<${container_t1} const> b_, ${INDEX_TYPE} _ldb, - ${DATA_TYPE} _beta, BufferIterator<${container_t2}> _C, ${INDEX_TYPE} _ldc, + ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, ${DATA_TYPE} _alpha, + BufferIterator<${DATA_TYPE} const> a_, ${INDEX_TYPE} _lda, + BufferIterator<${DATA_TYPE} const> b_, ${INDEX_TYPE} _ldb, + ${DATA_TYPE} _beta, BufferIterator<${DATA_TYPE}> _C, ${INDEX_TYPE} _ldc, const typename SB_Handle::event_t& _dependencies); 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, BufferIterator<${container_t0} const> a_, - ${INDEX_TYPE} _lda, BufferIterator<${container_t1} const> b_, ${INDEX_TYPE} _ldb, - ${DATA_TYPE} _beta, BufferIterator<${container_t2}> _C, ${INDEX_TYPE} _ldc, + ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, ${DATA_TYPE} _alpha, + BufferIterator<${DATA_TYPE} const> a_, ${INDEX_TYPE} _lda, + BufferIterator<${DATA_TYPE} const> b_, ${INDEX_TYPE} _ldb, + ${DATA_TYPE} _beta, BufferIterator<${DATA_TYPE}> _C, ${INDEX_TYPE} _ldc, ${INDEX_TYPE} batch_size, gemm_batch_type_t batch_type, const typename SB_Handle::event_t& _dependencies); template typename SB_Handle::event_t _gemm_strided_batched( SB_Handle& sb_handle, char _TransA, char _TransB, ${INDEX_TYPE} _M, - ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, ${DATA_TYPE} _alpha, BufferIterator<${container_t0} const> a_, - ${INDEX_TYPE} _lda, ${INDEX_TYPE} _stridea, BufferIterator<${container_t1} const> b_, + ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, ${DATA_TYPE} _alpha, + BufferIterator<${DATA_TYPE} const> a_, ${INDEX_TYPE} _lda, + ${INDEX_TYPE} _stridea, BufferIterator<${DATA_TYPE} const> b_, ${INDEX_TYPE} _ldb, ${INDEX_TYPE} _strideb, ${DATA_TYPE} _beta, - BufferIterator<${container_t2}> _C, ${INDEX_TYPE} _ldc, ${INDEX_TYPE} _stridec, - ${INDEX_TYPE} batch_size, const typename SB_Handle::event_t& _dependencies); + BufferIterator<${DATA_TYPE}> _C, ${INDEX_TYPE} _ldc, + ${INDEX_TYPE} _stridec, ${INDEX_TYPE} batch_size, + const typename SB_Handle::event_t& _dependencies); #endif } // namespace internal diff --git a/src/interface/blas3/gemm_launcher.cpp.in b/src/interface/blas3/gemm_launcher.cpp.in index 710376bc6..cecf37f2d 100644 --- a/src/interface/blas3/gemm_launcher.cpp.in +++ b/src/interface/blas3/gemm_launcher.cpp.in @@ -35,23 +35,10 @@ namespace blas { // BufferIterator -template class Gemm_Launcher< - BufferIterator<${container_t0}>, BufferIterator<${container_t1}>, - BufferIterator<${container_t2}>, ${WG_SIZE}, - ${DOUBLE_BUFFER}, ${CONFLICT_A}, ${CONFLICT_B}, ${CL_SIZE}, - Tile<${TIR}, ${TIC}, ${TWR}, ${TWC}, ${TSR}, ${TSC}, ${TLR}, ${TLC}, ${TIB}, - ${TWB}, ${JM_M}, ${JM_N}, ${JM_K}, ${JM_IN_T}, ${JM_OUT_T}>, - ${TRANS_A}, ${TRANS_B}, ${SYMM_A}, ${SYMM_B}, - static_cast(gemm_memory_t::${GEMM_MEMORY_TYPE}), - static_cast(gemm_algorithm_t::${GEMM_SHAPE_TYPE}), - static_cast(gemm_vectorization_t::${GEMM_VECTORIZE_TYPE}), - ${IS_BETA_ZERO}, ${VECTOR_SIZE}, - static_cast(gemm_batch_type_t::${BATCH_TYPE}), ${USE_JOINT_MATRIX}>; - template typename SB_Handle::event_t Gemm_Launcher< - BufferIterator<${container_t0}>, BufferIterator<${container_t1}>, - BufferIterator<${container_t2}>, ${WG_SIZE}, - ${DOUBLE_BUFFER}, ${CONFLICT_A}, ${CONFLICT_B}, ${CL_SIZE}, + BufferIterator<${DATA_TYPE}>, BufferIterator<${DATA_TYPE}>, + BufferIterator<${DATA_TYPE}>, ${WG_SIZE}, ${DOUBLE_BUFFER}, + ${CONFLICT_A}, ${CONFLICT_B}, ${CL_SIZE}, Tile<${TIR}, ${TIC}, ${TWR}, ${TWC}, ${TSR}, ${TSC}, ${TLR}, ${TLC}, ${TIB}, ${TWB}, ${JM_M}, ${JM_N}, ${JM_K}, ${JM_IN_T}, ${JM_OUT_T}>, ${TRANS_A}, ${TRANS_B}, ${SYMM_A}, ${SYMM_B}, @@ -63,30 +50,18 @@ template typename SB_Handle::event_t Gemm_Launcher< ${USE_JOINT_MATRIX}>::_select_gemm( SB_Handle& sb_handle, ${INDEX_TYPE} _M, ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, - ${DATA_TYPE} _alpha, BufferIterator<${container_t0}> a_, ${INDEX_TYPE} _lda, - ${INDEX_TYPE} _stridea, BufferIterator<${container_t1}> b_, ${INDEX_TYPE} _ldb, - ${INDEX_TYPE} _strideb, ${DATA_TYPE} _beta, BufferIterator<${container_t2}> _C, - ${INDEX_TYPE} _ldc, ${INDEX_TYPE} _stridec, ${INDEX_TYPE} batch_size, + ${DATA_TYPE} _alpha, BufferIterator<${DATA_TYPE}> a_, ${INDEX_TYPE} _lda, + ${INDEX_TYPE} _stridea, BufferIterator<${DATA_TYPE}> b_, + ${INDEX_TYPE} _ldb, ${INDEX_TYPE} _strideb, ${DATA_TYPE} _beta, + BufferIterator<${DATA_TYPE}> _C, ${INDEX_TYPE} _ldc, + ${INDEX_TYPE} _stridec, ${INDEX_TYPE} batch_size, const typename SB_Handle::event_t& _dependencies); #ifdef BLAS_ENABLE_CONST_INPUT -template class Gemm_Launcher< - BufferIterator<${container_t0} const>, BufferIterator<${container_t1} const>, - BufferIterator<${container_t2}>, ${WG_SIZE}, - ${DOUBLE_BUFFER}, ${CONFLICT_A}, ${CONFLICT_B}, ${CL_SIZE}, - Tile<${TIR}, ${TIC}, ${TWR}, ${TWC}, ${TSR}, ${TSC}, ${TLR}, ${TLC}, ${TIB}, - ${TWB}, ${JM_M}, ${JM_N}, ${JM_K}, ${JM_IN_T}, ${JM_OUT_T}>, - ${TRANS_A}, ${TRANS_B}, ${SYMM_A}, ${SYMM_B}, - static_cast(gemm_memory_t::${GEMM_MEMORY_TYPE}), - static_cast(gemm_algorithm_t::${GEMM_SHAPE_TYPE}), - static_cast(gemm_vectorization_t::${GEMM_VECTORIZE_TYPE}), - ${IS_BETA_ZERO}, ${VECTOR_SIZE}, - static_cast(gemm_batch_type_t::${BATCH_TYPE}), ${USE_JOINT_MATRIX}>; - template typename SB_Handle::event_t Gemm_Launcher< - BufferIterator<${container_t0} const>, BufferIterator<${container_t1} const>, - BufferIterator<${container_t2}>, ${WG_SIZE}, - ${DOUBLE_BUFFER}, ${CONFLICT_A}, ${CONFLICT_B}, ${CL_SIZE}, + BufferIterator<${DATA_TYPE} const>, BufferIterator<${DATA_TYPE} const >, + BufferIterator<${DATA_TYPE}>, ${WG_SIZE}, ${DOUBLE_BUFFER}, + ${CONFLICT_A}, ${CONFLICT_B}, ${CL_SIZE}, Tile<${TIR}, ${TIC}, ${TWR}, ${TWC}, ${TSR}, ${TSC}, ${TLR}, ${TLC}, ${TIB}, ${TWB}, ${JM_M}, ${JM_N}, ${JM_K}, ${JM_IN_T}, ${JM_OUT_T}>, ${TRANS_A}, ${TRANS_B}, ${SYMM_A}, ${SYMM_B}, @@ -95,32 +70,22 @@ template typename SB_Handle::event_t Gemm_Launcher< static_cast(gemm_vectorization_t::${GEMM_VECTORIZE_TYPE}), ${IS_BETA_ZERO}, ${VECTOR_SIZE}, static_cast(gemm_batch_type_t::${BATCH_TYPE}), - ${USE_JOINT_MATRIX}>::_select_gemm( - SB_Handle& sb_handle, ${INDEX_TYPE} _M, ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, - ${DATA_TYPE} _alpha, BufferIterator<${container_t0} const> a_, ${INDEX_TYPE} _lda, - ${INDEX_TYPE} _stridea, BufferIterator<${container_t1} const> b_, ${INDEX_TYPE} _ldb, - ${INDEX_TYPE} _strideb, ${DATA_TYPE} _beta, BufferIterator<${container_t2}> _C, - ${INDEX_TYPE} _ldc, ${INDEX_TYPE} _stridec, ${INDEX_TYPE} batch_size, - const typename SB_Handle::event_t& _dependencies); + ${USE_JOINT_MATRIX} > + ::_select_gemm( + SB_Handle& sb_handle, ${INDEX_TYPE} _M, ${INDEX_TYPE} _N, + ${INDEX_TYPE} _K, ${DATA_TYPE} _alpha, + BufferIterator<${DATA_TYPE} const> a_, ${INDEX_TYPE} _lda, + ${INDEX_TYPE} _stridea, BufferIterator<${DATA_TYPE} const> b_, + ${INDEX_TYPE} _ldb, ${INDEX_TYPE} _strideb, ${DATA_TYPE} _beta, + BufferIterator<${DATA_TYPE}> _C, ${INDEX_TYPE} _ldc, + ${INDEX_TYPE} _stridec, ${INDEX_TYPE} batch_size, + const typename SB_Handle::event_t& _dependencies); #endif #ifdef SB_ENABLE_USM // pointer -template class Gemm_Launcher< - ${container_t0}*, ${container_t1}*, ${container_t2}*, ${WG_SIZE}, - ${DOUBLE_BUFFER}, ${CONFLICT_A}, ${CONFLICT_B}, ${CL_SIZE}, - Tile<${TIR}, ${TIC}, ${TWR}, ${TWC}, ${TSR}, ${TSC}, ${TLR}, ${TLC}, ${TIB}, - ${TWB}, ${JM_M}, ${JM_N}, ${JM_K}, ${JM_IN_T}, ${JM_OUT_T}>, - ${TRANS_A}, ${TRANS_B}, ${SYMM_A}, ${SYMM_B}, - static_cast(gemm_memory_t::${GEMM_MEMORY_TYPE}), - static_cast(gemm_algorithm_t::${GEMM_SHAPE_TYPE}), - static_cast(gemm_vectorization_t::${GEMM_VECTORIZE_TYPE}), - ${IS_BETA_ZERO}, ${VECTOR_SIZE}, - static_cast(gemm_batch_type_t::${BATCH_TYPE}), ${USE_JOINT_MATRIX}>; - template typename SB_Handle::event_t Gemm_Launcher< - ${container_t0}*, ${container_t1}*, ${container_t2}*, ${WG_SIZE}, + ${DATA_TYPE}*, ${DATA_TYPE}*, ${DATA_TYPE}*, ${WG_SIZE}, ${DOUBLE_BUFFER}, ${CONFLICT_A}, ${CONFLICT_B}, ${CL_SIZE}, Tile<${TIR}, ${TIC}, ${TWR}, ${TWC}, ${TSR}, ${TSC}, ${TLR}, ${TLC}, ${TIB}, ${TWB}, ${JM_M}, ${JM_N}, ${JM_K}, ${JM_IN_T}, ${JM_OUT_T}>, @@ -133,28 +98,16 @@ template typename SB_Handle::event_t Gemm_Launcher< ${USE_JOINT_MATRIX}>::_select_gemm( SB_Handle& sb_handle, ${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, + ${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, const typename SB_Handle::event_t& _dependencies); // const pointer -template class Gemm_Launcher< - const ${container_t0}*, const ${container_t1}*, ${container_t2}*, ${WG_SIZE}, - ${DOUBLE_BUFFER}, ${CONFLICT_A}, ${CONFLICT_B}, ${CL_SIZE}, - Tile<${TIR}, ${TIC}, ${TWR}, ${TWC}, ${TSR}, ${TSC}, ${TLR}, ${TLC}, ${TIB}, - ${TWB}, ${JM_M}, ${JM_N}, ${JM_K}, ${JM_IN_T}, ${JM_OUT_T}>, - ${TRANS_A}, ${TRANS_B}, ${SYMM_A}, ${SYMM_B}, - static_cast(gemm_memory_t::${GEMM_MEMORY_TYPE}), - static_cast(gemm_algorithm_t::${GEMM_SHAPE_TYPE}), - static_cast(gemm_vectorization_t::${GEMM_VECTORIZE_TYPE}), - ${IS_BETA_ZERO}, ${VECTOR_SIZE}, - static_cast(gemm_batch_type_t::${BATCH_TYPE}), ${USE_JOINT_MATRIX}>; - template typename SB_Handle::event_t Gemm_Launcher< - const ${container_t0}*, const ${container_t1}*, ${container_t2}*, ${WG_SIZE}, - ${DOUBLE_BUFFER}, ${CONFLICT_A}, ${CONFLICT_B}, ${CL_SIZE}, + const ${DATA_TYPE}*, const ${DATA_TYPE}*, ${DATA_TYPE}*, + ${WG_SIZE}, ${DOUBLE_BUFFER}, ${CONFLICT_A}, ${CONFLICT_B}, ${CL_SIZE}, Tile<${TIR}, ${TIC}, ${TWR}, ${TWC}, ${TSR}, ${TSC}, ${TLR}, ${TLC}, ${TIB}, ${TWB}, ${JM_M}, ${JM_N}, ${JM_K}, ${JM_IN_T}, ${JM_OUT_T}>, ${TRANS_A}, ${TRANS_B}, ${SYMM_A}, ${SYMM_B}, @@ -166,9 +119,9 @@ template typename SB_Handle::event_t Gemm_Launcher< ${USE_JOINT_MATRIX}>::_select_gemm( SB_Handle& sb_handle, ${INDEX_TYPE} _M, ${INDEX_TYPE} _N, ${INDEX_TYPE} _K, - ${DATA_TYPE} _alpha, const ${container_t0}* a_, ${INDEX_TYPE} _lda, - ${INDEX_TYPE} _stridea, const ${container_t1}* b_, ${INDEX_TYPE} _ldb, - ${INDEX_TYPE} _strideb, ${DATA_TYPE} _beta, ${container_t2}* _C, + ${DATA_TYPE} _alpha, const ${DATA_TYPE} * a_, ${INDEX_TYPE} _lda, + ${INDEX_TYPE} _stridea, const ${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, const typename SB_Handle::event_t& _dependencies); #endif diff --git a/src/interface/blas3/symm.cpp.in b/src/interface/blas3/symm.cpp.in index 6a30d9e85..a4a38bfb9 100644 --- a/src/interface/blas3/symm.cpp.in +++ b/src/interface/blas3/symm.cpp.in @@ -32,24 +32,24 @@ namespace internal { template typename SB_Handle::event_t _symm( SB_Handle& sb_handle, char _side, char _uplo, ${INDEX_TYPE} _M, - ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, BufferIterator<${container_t0}> a_, - ${INDEX_TYPE} _lda, BufferIterator<${container_t1}> b_, ${INDEX_TYPE} _ldb, - ${DATA_TYPE} _beta, BufferIterator<${container_t2}> _C, ${INDEX_TYPE} _ldc, + ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, BufferIterator<${DATA_TYPE}> a_, + ${INDEX_TYPE} _lda, BufferIterator<${DATA_TYPE}> b_, ${INDEX_TYPE} _ldb, + ${DATA_TYPE} _beta, BufferIterator<${DATA_TYPE}> _C, ${INDEX_TYPE} _ldc, const typename SB_Handle::event_t& dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _symm( SB_Handle& sb_handle, char _side, char _uplo, ${INDEX_TYPE} _M, - ${INDEX_TYPE} _N, ${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, ${DATA_TYPE} _alpha, ${DATA_TYPE} * a_, + ${INDEX_TYPE} _lda, ${DATA_TYPE} * b_, ${INDEX_TYPE} _ldb, + ${DATA_TYPE} _beta, ${DATA_TYPE} * _C, ${INDEX_TYPE} _ldc, const typename SB_Handle::event_t& dependencies); template typename SB_Handle::event_t _symm( SB_Handle& sb_handle, char _side, char _uplo, ${INDEX_TYPE} _M, - ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, const ${container_t0}* a_, - ${INDEX_TYPE} _lda, const ${container_t1}* b_, ${INDEX_TYPE} _ldb, - ${DATA_TYPE} _beta, ${container_t2}* _C, ${INDEX_TYPE} _ldc, + ${INDEX_TYPE} _N, ${DATA_TYPE} _alpha, const ${DATA_TYPE} * a_, + ${INDEX_TYPE} _lda, const ${DATA_TYPE} * b_, ${INDEX_TYPE} _ldb, + ${DATA_TYPE} _beta, ${DATA_TYPE} * _C, ${INDEX_TYPE} _ldc, const typename SB_Handle::event_t& dependencies); #endif diff --git a/src/interface/blas3/trsm.cpp.in b/src/interface/blas3/trsm.cpp.in index dd30afd2c..ca9c6956e 100644 --- a/src/interface/blas3/trsm.cpp.in +++ b/src/interface/blas3/trsm.cpp.in @@ -33,22 +33,23 @@ namespace internal { template typename SB_Handle::event_t _trsm( SB_Handle& sb_handle, char side, char uplo, char trans, char diag, - ${INDEX_TYPE} M, ${INDEX_TYPE} N, ${DATA_TYPE} alpha, BufferIterator<${container_t0}> A, - ${INDEX_TYPE} lda, BufferIterator<${container_t1}> B, ${INDEX_TYPE} ldb, + ${INDEX_TYPE} M, ${INDEX_TYPE} N, ${DATA_TYPE} alpha, + BufferIterator<${DATA_TYPE}> A, ${INDEX_TYPE} lda, + BufferIterator<${DATA_TYPE}> B, ${INDEX_TYPE} ldb, const typename SB_Handle::event_t& _dependencies); #ifdef SB_ENABLE_USM template typename SB_Handle::event_t _trsm( SB_Handle& sb_handle, char side, char uplo, char trans, char diag, - ${INDEX_TYPE} M, ${INDEX_TYPE} N, ${DATA_TYPE} alpha, ${container_t0}* A, - ${INDEX_TYPE} lda, ${container_t1}* B, ${INDEX_TYPE} ldb, + ${INDEX_TYPE} M, ${INDEX_TYPE} N, ${DATA_TYPE} alpha, ${DATA_TYPE} * A, + ${INDEX_TYPE} lda, ${DATA_TYPE} * B, ${INDEX_TYPE} ldb, const typename SB_Handle::event_t& _dependencies); template typename SB_Handle::event_t _trsm( SB_Handle& sb_handle, char side, char uplo, char trans, char diag, - ${INDEX_TYPE} M, ${INDEX_TYPE} N, ${DATA_TYPE} alpha, const ${container_t0}* A, - ${INDEX_TYPE} lda, ${container_t1}* B, ${INDEX_TYPE} ldb, - const typename SB_Handle::event_t& _dependencies); + ${INDEX_TYPE} M, ${INDEX_TYPE} N, ${DATA_TYPE} alpha, + const ${DATA_TYPE} * A, ${INDEX_TYPE} lda, ${DATA_TYPE} * B, + ${INDEX_TYPE} ldb, const typename SB_Handle::event_t& _dependencies); #endif } // namespace internal diff --git a/src/interface/reduction/reduction.cpp.in b/src/interface/reduction/reduction.cpp.in index 56470eade..9562db02b 100644 --- a/src/interface/reduction/reduction.cpp.in +++ b/src/interface/reduction/reduction.cpp.in @@ -23,38 +23,38 @@ * **************************************************************************/ -#include "operations/extension/reduction.hpp" #include "container/sycl_iterator.hpp" -#include "sb_handle/sycl_blas_handle.hpp" -#include "sb_handle/kernel_constructor.hpp" #include "interface/blas1_interface.hpp" #include "interface/reduction_interface.hpp" #include "operations/blas_constants.hpp" +#include "operations/extension/reduction.hpp" +#include "sb_handle/kernel_constructor.hpp" +#include "sb_handle/sycl_blas_handle.hpp" #include "views/view_sycl.hpp" namespace blas { namespace extension { namespace internal { -template -typename SB_Handle::event_t _reduction<${OPERATOR}, ${DATA_TYPE}>( - SB_Handle& sb_handle, BufferIterator<${container_t0}> buffer_in, ${INDEX_TYPE} ld, - BufferIterator<${container_t1}> buffer_out, ${INDEX_TYPE} rows, ${INDEX_TYPE} cols, reduction_dim_t reduction_type, +template typename SB_Handle::event_t _reduction<${OPERATOR}, ${DATA_TYPE}>( + SB_Handle& sb_handle, BufferIterator<${DATA_TYPE}> buffer_in, + ${INDEX_TYPE} ld, BufferIterator<${DATA_TYPE}> buffer_out, + ${INDEX_TYPE} rows, ${INDEX_TYPE} cols, reduction_dim_t reduction_type, const typename SB_Handle::event_t& dependencies); #ifdef BLAS_ENABLE_CONST_INPUT -template - typename SB_Handle::event_t _reduction<${OPERATOR}, ${DATA_TYPE}>( - SB_Handle& sb_handle, BufferIterator<${container_t0} const> buffer_in, ${INDEX_TYPE} ld, - BufferIterator<${container_t1}> buffer_out, ${INDEX_TYPE} rows, ${INDEX_TYPE} cols, reduction_dim_t reduction_type, - const typename SB_Handle::event_t& dependencies); +template typename SB_Handle::event_t _reduction<${OPERATOR}, ${DATA_TYPE}>( + SB_Handle& sb_handle, BufferIterator<${DATA_TYPE} const> buffer_in, + ${INDEX_TYPE} ld, BufferIterator<${DATA_TYPE}> buffer_out, + ${INDEX_TYPE} rows, ${INDEX_TYPE} cols, reduction_dim_t reduction_type, + const typename SB_Handle::event_t& dependencies); #endif #ifdef SB_ENABLE_USM -template -typename SB_Handle::event_t _reduction<${OPERATOR}, ${DATA_TYPE}>( - SB_Handle& sb_handle, ${container_t0}* buffer_in, ${INDEX_TYPE} ld, - ${container_t1}* buffer_out, ${INDEX_TYPE} rows, ${INDEX_TYPE} cols, reduction_dim_t reduction_type, +template typename SB_Handle::event_t _reduction<${OPERATOR}, ${DATA_TYPE}>( + SB_Handle& sb_handle, ${DATA_TYPE} * buffer_in, ${INDEX_TYPE} ld, + ${DATA_TYPE} * buffer_out, ${INDEX_TYPE} rows, ${INDEX_TYPE} cols, + reduction_dim_t reduction_type, const typename SB_Handle::event_t& dependencies); #endif diff --git a/src/operations/blas2/gemv.hpp b/src/operations/blas2/gemv.hpp index 43d15b8cf..c3d563bc8 100644 --- a/src/operations/blas2/gemv.hpp +++ b/src/operations/blas2/gemv.hpp @@ -53,7 +53,7 @@ SYCL_BLAS_INLINE bool SumMatrixColumns::valid_thread( template SYCL_BLAS_INLINE typename SumMatrixColumns::value_t -SumMatrixColumns::eval(typename SumMatrixColumns::index_t i) const { +SumMatrixColumns::eval(typename SumMatrixColumns::index_t i) { auto dimR = rhs_.get_size_row(); auto dimC = rhs_.get_size_col();