Skip to content

Commit d18f338

Browse files
devrealdevreal
committed
GEMM: Add a dense version using either TT or make_tt
Signed-off-by: Joseph Schuchart <joseph.schuchart@stonybrook.edu>
1 parent d07e51e commit d18f338

5 files changed

Lines changed: 2249 additions & 0 deletions

File tree

‎examples/CMakeLists.txt‎

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -15,6 +15,12 @@ if (TARGET tiledarray)
1515
add_ttg_executable(bspmm spmm/spmm.cc LINK_LIBRARIES tiledarray TiledArray_Eigen BTAS
1616
COMPILE_DEFINITIONS BLOCK_SPARSE_GEMM=1;BTAS_TARGET_MAX_INDEX_RANK=2)
1717

18+
# Dense 2.5D SUMMA (no sparsity handling)
19+
add_ttg_executable(gemm gemm/gemm.cc LINK_LIBRARIES tiledarray BTAS
20+
COMPILE_DEFINITIONS BTAS_TARGET_MAX_INDEX_RANK=2)
21+
add_ttg_executable(gemm2 gemm/gemm2.cc LINK_LIBRARIES tiledarray BTAS
22+
COMPILE_DEFINITIONS BTAS_TARGET_MAX_INDEX_RANK=2)
23+
1824
add_ttg_executable(testing_dpotrf potrf/testing_dpotrf.cc LINK_LIBRARIES tiledarray lapackpp)
1925
add_ttg_executable(testing_dtrtri potrf/testing_dtrtri.cc LINK_LIBRARIES tiledarray lapackpp)
2026
add_ttg_executable(testing_dlauum potrf/testing_dlauum.cc LINK_LIBRARIES tiledarray lapackpp)
@@ -25,6 +31,14 @@ if (TARGET tiledarray)
2531
LINK_LIBRARIES tiledarray TiledArray_Eigen BTAS CUDA::cublas
2632
COMPILE_DEFINITIONS BLOCK_SPARSE_GEMM=1;BTAS_TARGET_MAX_INDEX_RANK=2;TTG_ENABLE_CUDA=1
2733
RUNTIMES "parsec")
34+
add_ttg_executable(gemm-cuda gemm/gemm.cc
35+
LINK_LIBRARIES tiledarray BTAS CUDA::cublas
36+
COMPILE_DEFINITIONS BTAS_TARGET_MAX_INDEX_RANK=2;TTG_ENABLE_CUDA=1
37+
RUNTIMES "parsec")
38+
add_ttg_executable(gemm2-cuda gemm/gemm2.cc
39+
LINK_LIBRARIES tiledarray BTAS CUDA::cublas
40+
COMPILE_DEFINITIONS BTAS_TARGET_MAX_INDEX_RANK=2;TTG_ENABLE_CUDA=1
41+
RUNTIMES "parsec")
2842

2943
if (TARGET CUDA::cusolver)
3044
add_ttg_executable(testing_dpotrf_cuda potrf/testing_dpotrf.cc
@@ -37,6 +51,14 @@ if (TARGET tiledarray)
3751
LINK_LIBRARIES tiledarray TiledArray_Eigen roc::hipblas
3852
COMPILE_DEFINITIONS BLOCK_SPARSE_GEMM=1;BTAS_TARGET_MAX_INDEX_RANK=2;TTG_ENABLE_HIP=1
3953
RUNTIMES "parsec")
54+
add_ttg_executable(gemm-hip gemm/gemm.cc
55+
LINK_LIBRARIES tiledarray roc::hipblas
56+
COMPILE_DEFINITIONS BTAS_TARGET_MAX_INDEX_RANK=2;TTG_ENABLE_HIP=1
57+
RUNTIMES "parsec")
58+
add_ttg_executable(gemm2-hip gemm/gemm2.cc
59+
LINK_LIBRARIES tiledarray roc::hipblas
60+
COMPILE_DEFINITIONS BTAS_TARGET_MAX_INDEX_RANK=2;TTG_ENABLE_HIP=1
61+
RUNTIMES "parsec")
4062
if (TARGET roc::hipsolver)
4163
add_ttg_executable(testing_dpotrf_hip potrf/testing_dpotrf.cc
4264
LINK_LIBRARIES lapackpp tiledarray roc::hipblas roc::hipsolver
@@ -48,6 +70,14 @@ if (TARGET tiledarray)
4870
LINK_LIBRARIES tiledarray TiledArray_Eigen BTAS MKL::MKL_DPCPP level_zero::ze_loader m
4971
COMPILE_DEFINITIONS BLOCK_SPARSE_GEMM=1;BTAS_TARGET_MAX_INDEX_RANK=2;TTG_ENABLE_LEVEL_ZERO=1
5072
RUNTIMES "parsec")
73+
add_ttg_executable(gemm-lz gemm/gemm.cc
74+
LINK_LIBRARIES tiledarray BTAS MKL::MKL_DPCPP level_zero::ze_loader m
75+
COMPILE_DEFINITIONS BTAS_TARGET_MAX_INDEX_RANK=2;TTG_ENABLE_LEVEL_ZERO=1
76+
RUNTIMES "parsec")
77+
add_ttg_executable(gemm2-lz gemm/gemm2.cc
78+
LINK_LIBRARIES tiledarray BTAS MKL::MKL_DPCPP level_zero::ze_loader m
79+
COMPILE_DEFINITIONS BTAS_TARGET_MAX_INDEX_RANK=2;TTG_ENABLE_LEVEL_ZERO=1
80+
RUNTIMES "parsec")
5181
endif()
5282

5383
if (TTG_HAVE_CUDA)

‎examples/gemm/devicegemm.h‎

Lines changed: 82 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,82 @@
1+
// SPDX-License-Identifier: BSD-3-Clause
2+
3+
#if defined(TTG_ENABLE_LEVEL_ZERO)
4+
#include <oneapi/mkl.hpp>
5+
#include <sys/time.h>
6+
#endif
7+
8+
#include "../devblas_helper.h"
9+
10+
11+
template <typename Blk>
12+
inline void device_gemm(Blk &C, const Blk &A, const Blk &B) {
13+
using blk_t = Blk;
14+
using T = typename blk_t::value_type;
15+
static_assert(std::is_same_v<T,double> || std::is_same_v<T,float>);
16+
static const T alpha = 1.0;
17+
static const T beta = 1.0;
18+
// make sure all memory is on the device
19+
// TODO: A and B are read-only so the owner device will be 0. How to fix?
20+
//assert(A.b.get_current_device() != 0);
21+
//assert(B.b.get_current_device() != 0);
22+
auto device = ttg::device::current_device();
23+
assert(device.is_device());
24+
#if defined(TTG_ENABLE_CUDA)
25+
if constexpr (std::is_same_v<T,double>) {
26+
cublasDgemm(cublas_handle(), CUBLAS_OP_N, CUBLAS_OP_N, C.extent(0), C.extent(1), A.extent(1),
27+
&alpha, A.b.current_device_ptr(), A.extent(0), B.b.current_device_ptr(), B.extent(0), &beta,
28+
C.b.current_device_ptr(), C.extent(0));
29+
}
30+
else if constexpr (std::is_same_v<T,float>) {
31+
cublasSgemm(cublas_handle(), CUBLAS_OP_N, CUBLAS_OP_N, C.extent(0), C.extent(1), A.extent(1),
32+
&alpha, A.b.current_device_ptr(), A.extent(0), B.b.current_device_ptr(), B.extent(0), &beta,
33+
C.b.current_device_ptr(), C.extent(0));
34+
}
35+
#elif defined(TTG_ENABLE_HIP)
36+
if constexpr (std::is_same_v<T,double>) {
37+
hipblasDgemm(hipblas_handle(),
38+
HIPBLAS_OP_N, HIPBLAS_OP_N,
39+
C.extent(0), C.extent(1), A.extent(1), &alpha,
40+
A.b.current_device_ptr(), A.extent(0),
41+
B.b.current_device_ptr(), B.extent(0), &beta,
42+
C.b.current_device_ptr(), C.extent(0));
43+
} else if constexpr (std::is_same_v<T,float>) {
44+
hipblasSgemm(hipblas_handle(),
45+
HIPBLAS_OP_N, HIPBLAS_OP_N,
46+
C.extent(0), C.extent(1), A.extent(1), &alpha,
47+
A.b.current_device_ptr(), A.extent(0),
48+
B.b.current_device_ptr(), B.extent(0), &beta,
49+
C.b.current_device_ptr(), C.extent(0));
50+
}
51+
#elif defined(TTG_ENABLE_LEVEL_ZERO)
52+
53+
#if defined(DEBUG_SYNCHRONOUS)
54+
try {
55+
#endif /* DEBUG_SYNCHRONOUS */
56+
cl::sycl::event gemm_event;
57+
gemm_event = oneapi::mkl::blas::gemm(ttg::device::current_stream(),
58+
oneapi::mkl::transpose::N, oneapi::mkl::transpose::N,
59+
C.extent(0), C.extent(1), A.extent(1),
60+
alpha, A.b.current_device_ptr(), A.extent(0),
61+
B.b.current_device_ptr(), B.extent(0),
62+
beta, C.b.current_device_ptr(), C.extent(0));
63+
#if defined(DEBUG_SYNCHRONOUS)
64+
gemm_event.wait();
65+
} catch (const oneapi::mkl::invalid_argument &e) {
66+
std::cerr << "OneAPI MKL BLAS GEMM throws invalid argument exception" << std::endl;
67+
} catch (const oneapi::mkl::unsupported_device &e) {
68+
std::cerr << "OneAPI MKL BLAS GEMM throws unsuported device exception" << std::endl;
69+
} catch (const oneapi::mkl::host_bad_alloc &e) {
70+
std::cerr << "OneAPI MKL BLAS GEMM throws host bad allocation exception" << std::endl;
71+
} catch (const oneapi::mkl::device_bad_alloc &e) {
72+
std::cerr << "OneAPI MKL BLAS GEMM throws device bad allocation exception" << std::endl;
73+
} catch (const oneapi::mkl::unimplemented &e) {
74+
std::cerr << "OneAPI MKL BLAS GEMM throws unimplemented exception" << std::endl;
75+
} catch (const std::exception& e) {
76+
std::cerr << "OneAPI MKL BLAS GEMM throws unexpected exception" << std::endl;
77+
} catch (...) {
78+
std::cerr << "OneAPI MKL BLAS GEMM throws unexpected exception that is also badly formatted..." << std::endl;
79+
}
80+
#endif /* DEBUG_SYNCHRONOUS */
81+
#endif
82+
}

0 commit comments

Comments
 (0)