ContextQMD
Back to Cutlass

CUTLASS: matrix.h Source File

docs/layout_2matrix_8h_source.html

4.4.289.7 KB
Original Source

| | CUTLASS

CUDA Templates for Linear Algebra Subroutines and Solvers |

layout/matrix.h

Go to the documentation of this file.

1 /***************************************************************************************************

2 * Copyright (c) 2017-2019, NVIDIA CORPORATION. All rights reserved.

3 *

4 * Redistribution and use in source and binary forms, with or without modification, are permitted

5 * provided that the following conditions are met:

6 * * Redistributions of source code must retain the above copyright notice, this list of

7 * conditions and the following disclaimer.

8 * * Redistributions in binary form must reproduce the above copyright notice, this list of

9 * conditions and the following disclaimer in the documentation and/or other materials

10 * provided with the distribution.

11 * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used

12 * to endorse or promote products derived from this software without specific prior written

13 * permission.

14 *

15 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR

16 * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND

17 * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE

18 * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,

19 * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;

20 * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,

21 * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE

22 * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

23 *

24 **************************************************************************************************/

34 #pragma once

35

36 #include "cutlass/cutlass.h"

37 #include "cutlass/matrix_coord.h"

38 #include "cutlass/matrix_traits.h"

39

40 namespace cutlass {

41 namespace layout {

42

44 //

45 // Defines data layouts of various matrix formats usable by TensorRef and other classes.

46 //

48

50 class RowMajor {

51 public:

53static int const kRank = 2;

54

56static int const kStrideRank = 1;

57

59using Index = int32_t;

60

62using LongIndex = int64_t;

63

65using TensorCoord = MatrixCoord;

66

68using Stride = Coord<kStrideRank, Index>;

69

70 private:

71//

72// Data members

73//

74

76Stride stride_;

77

78 public:

79//

80// Methods

81//

82

84CUTLASS_HOST_DEVICE

85RowMajor(Index ldm = 0): stride_(ldm) { }

86

88CUTLASS_HOST_DEVICE

89RowMajor(Stride stride): stride_(stride) { }

90

92CUTLASS_HOST_DEVICE

93static RowMajor packed(MatrixCoord const &extent) {

94return RowMajor(extent.column());

95 }

96

99CUTLASS_HOST_DEVICE

100LongIndex operator()(MatrixCoord const &coord) const {

101return coord.row() * stride_[0] + coord.column();

102 }

103

105CUTLASS_HOST_DEVICE

106MatrixCoord inverse(LongIndex offset) const {

107return MatrixCoord(Index(offset / stride_[0]), Index(offset % stride_[0]));

108 }

109

111CUTLASS_HOST_DEVICE

112Stride stride() const {

113return stride_;

114 }

115

117CUTLASS_HOST_DEVICE

118Stride & stride() {

119return stride_;

120 }

121

123CUTLASS_HOST_DEVICE

124Index stride(int idx) const {

125return stride_[idx];

126 }

127

129CUTLASS_HOST_DEVICE

130Index & stride(int idx) {

131return stride_[idx];

132 }

133

135CUTLASS_HOST_DEVICE

136LongIndex capacity(MatrixCoord const &extent) const {

137return extent.row() * stride_[0];

138 }

139 };

140

142 class ColumnMajor {

143 public:

145static int const kRank = 2;

146

148static int const kStrideRank = 1;

149

151using Index = int32_t;

152

154using LongIndex = int64_t;

155

157using TensorCoord = MatrixCoord;

158

160using Stride = Coord<kStrideRank, Index>;

161

162 private:

163//

164// Data members

165//

166

168Stride stride_;

169

170 public:

171//

172// Methods

173//

174

176CUTLASS_HOST_DEVICE

177ColumnMajor(Index ldm = 0): stride_(ldm) { }

178

180CUTLASS_HOST_DEVICE

181ColumnMajor(Stride stride): stride_(stride) { }

182

183

185CUTLASS_HOST_DEVICE

186static ColumnMajor packed(MatrixCoord const &extent) {

187return ColumnMajor(extent.row());

188 }

189

192CUTLASS_HOST_DEVICE

193LongIndex operator()(MatrixCoord const &coord) const {

194return coord.row() + coord.column() * stride_[0];

195 }

196

198CUTLASS_HOST_DEVICE

199MatrixCoord inverse(LongIndex offset) const {

200return MatrixCoord(Index(offset % stride_[0]), Index(offset / stride_[0]));

201 }

202

204CUTLASS_HOST_DEVICE

205Stride stride() const {

206return stride_;

207 }

208

210CUTLASS_HOST_DEVICE

211Stride & stride() {

212return stride_;

213 }

214

216CUTLASS_HOST_DEVICE

217Index stride(int idx) const {

218return stride_[idx];

219 }

220

222CUTLASS_HOST_DEVICE

223Index & stride(int idx) {

224return stride_[idx];

225 }

226

228CUTLASS_HOST_DEVICE

229LongIndex capacity(MatrixCoord const &extent) const {

230return extent.column() * stride_[0];

231 }

232 };

233

236 template <int Interleave>

237 struct RowMajorInterleaved {

238

240static int const kRank = 2;

241

243static int const kStrideRank = 1;

244

246using Index = int32_t;

247

249using LongIndex = int64_t;

250

252using TensorCoord = MatrixCoord;

253

255using Stride = Coord<kStrideRank, Index>;

256

258static int const kInterleave = Interleave;

259

260 private:

261//

262// Data members

263//

264

266Stride stride_;

267

268 public:

269//

270// Methods

271//

272

274CUTLASS_HOST_DEVICE

275RowMajorInterleaved(Index ldm = 0): stride_(ldm) { }

276

278CUTLASS_HOST_DEVICE

279RowMajorInterleaved(Stride stride): stride_(stride) { }

280

282CUTLASS_HOST_DEVICE

283static RowMajorInterleaved packed(MatrixCoord const &extent) {

284return RowMajorInterleaved(extent.column() * kInterleave);

285 }

286

289CUTLASS_HOST_DEVICE

290LongIndex operator()(MatrixCoord const &coord) const {

291Index row_major = coord.row() / kInterleave;

292Index row_minor = coord.row() % kInterleave;

293return row_major * stride_[0] + coord.column() * kInterleave + row_minor;

294 }

295

297CUTLASS_HOST_DEVICE

298MatrixCoord inverse(LongIndex offset) const {

299

300Index row_major = Index(offset / stride_[0]);

301Index residual = Index(offset % stride_[0]);

302

303Index column = residual / kInterleave;

304Index row_minor = residual % kInterleave;

305

306return MatrixCoord(row_major * kInterleave + row_minor, column);

307 }

308

310CUTLASS_HOST_DEVICE

311Stride stride() const {

312return stride_;

313 }

314

316CUTLASS_HOST_DEVICE

317Stride & stride() {

318return stride_;

319 }

320

322CUTLASS_HOST_DEVICE

323Index stride(int idx) const {

324return stride_[idx];

325 }

326

328CUTLASS_HOST_DEVICE

329Index & stride(int idx) {

330return stride_[idx];

331 }

332

334CUTLASS_HOST_DEVICE

335LongIndex capacity(MatrixCoord const &extent) const {

336return (extent.row() + kInterleave - 1) / kInterleave * stride_[0];

337 }

338 };

339

342 template <int Interleave>

343 struct ColumnMajorInterleaved {

344

346static int const kRank = 2;

347

349static int const kStrideRank = 1;

350

352using Index = int32_t;

353

355using LongIndex = int64_t;

356

358using TensorCoord = MatrixCoord;

359

361using Stride = Coord<kStrideRank, Index>;

362

364static int const kInterleave = Interleave;

365

366 private:

367//

368// Data members

369//

370

372Stride stride_;

373

374 public:

375//

376// Methods

377//

378

380CUTLASS_HOST_DEVICE

381ColumnMajorInterleaved(Index ldm = 0): stride_(ldm) { }

382

384CUTLASS_HOST_DEVICE

385ColumnMajorInterleaved(Stride stride): stride_(stride) { }

386

387

389CUTLASS_HOST_DEVICE

390static ColumnMajorInterleaved packed(MatrixCoord const &extent) {

391return ColumnMajorInterleaved(extent.row() * kInterleave);

392 }

393

396CUTLASS_HOST_DEVICE

397LongIndex operator()(MatrixCoord const &coord) const {

398Index column_major = coord.column() / kInterleave;

399Index column_minor = coord.column() % kInterleave;

400return column_major * stride_[0] + coord.row() * kInterleave + column_minor;

401 }

402

404CUTLASS_HOST_DEVICE

405MatrixCoord inverse(LongIndex offset) const {

406

407Index column_major = Index(offset / stride_[0]);

408Index residual = Index(offset % stride_[0]);

409

410Index row = residual / kInterleave;

411Index column_minor = residual % kInterleave;

412

413return MatrixCoord(row, column_major * kInterleave + column_minor);

414 }

415

417CUTLASS_HOST_DEVICE

418Stride stride() const {

419return stride_;

420 }

421

423CUTLASS_HOST_DEVICE

424Stride & stride() {

425return stride_;

426 }

427

429CUTLASS_HOST_DEVICE

430Index stride(int idx) const {

431return stride_[idx];

432 }

433

435CUTLASS_HOST_DEVICE

436Index & stride(int idx) {

437return stride_[idx];

438 }

439

441CUTLASS_HOST_DEVICE

442LongIndex capacity(MatrixCoord const &extent) const {

443return (extent.column() + kInterleave - 1) / kInterleave * stride_[0];

444 }

445 };

446

448 enum class Matrix {

449kColumnMajor,

450kRowMajor

451 };

452

455 struct ContiguousMatrix {

456

458static int const kRank = 2;

459

461static int const kStrideRank = 1;

462

464using Index = int32_t;

465

467using LongIndex = int64_t;

468

470using TensorCoord = MatrixCoord;

471

473using Stride = Coord<kStrideRank, Index>;

474

475 private:

476//

477// Data members

478//

479

481Stride stride_;

482

484Matrix layout_;

485

486 public:

487//

488// Methods

489//

490

492CUTLASS_HOST_DEVICE

493ContiguousMatrix(

494Index ldm = 0,

495Matrix layout = Matrix::kColumnMajor

496 ):

497 stride_(ldm), layout_(layout) { }

498

500CUTLASS_HOST_DEVICE

501static ContiguousMatrix packed(

502MatrixCoord const &extent,

503Matrix layout = Matrix::kColumnMajor) {

504

505Index ldm = 0;

506if (layout == Matrix::kColumnMajor) {

507 ldm = extent.row();

508 }

509else if (layout == Matrix::kRowMajor) {

510 ldm = extent.column();

511 }

512return ContiguousMatrix(ldm, layout);

513 }

514

517CUTLASS_HOST_DEVICE

518LongIndex operator()(MatrixCoord const &coord) const {

519if (layout_ == Matrix::kColumnMajor) {

520return coord.row() + coord.column() * stride_[0];

521 }

522else if (layout_ == Matrix::kRowMajor) {

523return coord.row() * stride_[0] + coord.column();

524 }

525else {

526// degenerate case

527return 0;

528 }

529 }

530

532CUTLASS_HOST_DEVICE

533MatrixCoord inverse(LongIndex offset) const {

534// TODO

535return MatrixCoord(0, 0);

536 }

537

539CUTLASS_HOST_DEVICE

540Stride stride() const {

541return stride_;

542 }

543

545CUTLASS_HOST_DEVICE

546Stride & stride() {

547return stride_;

548 }

549

551CUTLASS_HOST_DEVICE

552Index stride(int idx) const {

553return stride_[idx];

554 }

555

557CUTLASS_HOST_DEVICE

558Index & stride(int idx) {

559return stride_[idx];

560 }

561

563CUTLASS_HOST_DEVICE

564LongIndex capacity(MatrixCoord const &extent) const {

565if (layout_ == Matrix::kColumnMajor) {

566return stride_[0] * extent.column();

567 }

568else if (layout_ == Matrix::kRowMajor) {

569return stride_[0] * extent.row();

570 }

571else {

572// degenerate case

573return 0;

574 }

575 }

576 };

577

580 template <int BlockRows, int BlockColumns>

581 struct ColumnMajorBlockLinear {

583static int const kRank = 2;

584

586static int const kStrideRank = 1;

587

589using Index = int32_t;

590

592using LongIndex = int64_t;

593

595using TensorCoord = MatrixCoord;

596

598using Stride = Coord<kStrideRank, Index>;

599

601static int const kBlockRows = BlockRows;

602

604static int const kBlockColumns = BlockColumns;

605

606 private:

607//

608// Data members

609//

610

612Stride stride_;

613

614 public:

615//

616// Methods

617//

618

620CUTLASS_HOST_DEVICE

621ColumnMajorBlockLinear(Index ldm = 0): stride_(ldm) { }

622

624CUTLASS_HOST_DEVICE

625static ColumnMajorBlockLinear packed(MatrixCoord const &extent) {

626return ColumnMajorBlockLinear(extent.row() * kBlockRows * kBlockColumns);

627 }

628

631CUTLASS_HOST_DEVICE

632LongIndex operator()(MatrixCoord const &coord) const {

633return

634 (coord.row() % kBlockRows) +

635 (coord.column() % kBlockColumns) * kBlockRows +

636 (coord.row() / kBlockRows) * kBlockRows * kBlockColumns +

637 (coord.column() / kBlockColumns) * stride_[0];

638 }

639

641CUTLASS_HOST_DEVICE

642MatrixCoord inverse(LongIndex offset) const {

643

644// TODO

645return MatrixCoord(0, 0);

646 }

647

649CUTLASS_HOST_DEVICE

650Stride stride() const {

651return stride_;

652 }

653

655CUTLASS_HOST_DEVICE

656Stride & stride() {

657return stride_;

658 }

659

661CUTLASS_HOST_DEVICE

662Index stride(int idx) const {

663return stride_[idx];

664 }

665

667CUTLASS_HOST_DEVICE

668Index & stride(int idx) {

669return stride_[idx];

670 }

671

673CUTLASS_HOST_DEVICE

674LongIndex capacity(MatrixCoord const &extent) const {

675return (extent.column() + kBlockColumns - 1) / kBlockColumns * stride_[0];

676 }

677 };

678

681 template <int BlockRows, int BlockColumns>

682 struct RowMajorBlockLinear {

684static int const kRank = 2;

685

687static int const kStrideRank = 1;

688

690using Index = int32_t;

691

693using LongIndex = int64_t;

694

696using TensorCoord = MatrixCoord;

697

699using Stride = Coord<kStrideRank, Index>;

700

702static int const kBlockRows = BlockRows;

703

705static int const kBlockColumns = BlockColumns;

706

707 private:

708//

709// Data members

710//

711

713Stride stride_;

714

715 public:

716//

717// Methods

718//

719

721CUTLASS_HOST_DEVICE

722RowMajorBlockLinear(Index ldm = 0): stride_(ldm) { }

723

725CUTLASS_HOST_DEVICE

726static RowMajorBlockLinear packed(MatrixCoord const &extent) {

727return RowMajorBlockLinear(extent.column() * kBlockRows * kBlockColumns);

728 }

729

732CUTLASS_HOST_DEVICE

733LongIndex operator()(MatrixCoord const &coord) const {

734return

735 (coord.column() % kBlockColumns) +

736 (coord.row() % kBlockRows) * kBlockColumns +

737 (coord.column() / kBlockColumns) * kBlockRows * kBlockColumns +

738 (coord.row() / kBlockRows) * stride_[0];

739 }

740

742CUTLASS_HOST_DEVICE

743MatrixCoord inverse(LongIndex offset) const {

744// TODO

745return MatrixCoord(0, 0);

746 }

747

749CUTLASS_HOST_DEVICE

750Stride stride() const {

751return stride_;

752 }

753

755CUTLASS_HOST_DEVICE

756Stride & stride() {

757return stride_;

758 }

759

761CUTLASS_HOST_DEVICE

762Index stride(int idx) const {

763return stride_[idx];

764 }

765

767CUTLASS_HOST_DEVICE

768Index & stride(int idx) {

769return stride_[idx];

770 }

771

773CUTLASS_HOST_DEVICE

774LongIndex capacity(MatrixCoord const &extent) const {

775return (extent.row() + kBlockRows - 1) / kBlockRows * stride_[0];

776 }

777 };

778

780

781 struct GeneralMatrix {

782

784static int const kRank = 2;

785

787static int const kStrideRank = 2;

788

790using Index = int32_t;

791

793using LongIndex = int64_t;

794

796using TensorCoord = MatrixCoord;

797

799using Stride = Coord<kStrideRank, Index>;

800

801 private:

802//

803// Data members

804//

805

806MatrixLayout layout_id_;

807

809Stride stride_;

810

811 public:

812//

813// Methods

814//

815

817CUTLASS_HOST_DEVICE

818GeneralMatrix(): layout_id_(MatrixLayout::kColumnMajor), stride_(make_Coord(0, 1)) { }

819

821CUTLASS_HOST_DEVICE

822GeneralMatrix(

823MatrixLayout layout_id,

824Index ldm,

825Index interleave): layout_id_(layout_id), stride_(make_Coord(ldm, interleave)) { }

826

828CUTLASS_HOST_DEVICE

829static GeneralMatrix packed(

830MatrixCoord const &extent,

831MatrixLayout layout_id = MatrixLayout::kColumnMajor,

832Index interleave = 1) {

833

834Index c;

835if (layout_id == MatrixLayout::kRowMajor) {

836 c = extent.column();

837 }

838else {

839 c = extent.row();

840 }

841

842Index ldm = c * interleave;

843

844return GeneralMatrix(layout_id, ldm, interleave);

845 }

846

849CUTLASS_HOST_DEVICE

850LongIndex operator()(MatrixCoord const &coord) const {

851Index c, s;

852if (layout_id_ == MatrixLayout::kRowMajor) {

853 c = coord.column();

854 s = coord.row();

855 }

856else {

857 s = coord.column();

858 c = coord.row();

859 }

860

861Index v = s / stride_[1];

862Index residual = (s % stride_[1]);

863

864return LongIndex(c) * LongIndex(stride_[1]) + LongIndex(v) * LongIndex(stride_[0]) + residual;

865 }

866

868CUTLASS_HOST_DEVICE

869Stride stride() const {

870return stride_;

871 }

872

873CUTLASS_HOST_DEVICE

874MatrixLayout layout_id() const {

875return layout_id_;

876 }

877

879CUTLASS_HOST_DEVICE

880Stride & stride() {

881return stride_;

882 }

883

884CUTLASS_HOST_DEVICE

885MatrixLayout & layout_id() {

886return layout_id_;

887 }

888

890CUTLASS_HOST_DEVICE

891Index stride(int idx) const {

892return stride_[idx];

893 }

894

896CUTLASS_HOST_DEVICE

897Index & stride(int idx) {

898return stride_[idx];

899 }

900

902CUTLASS_HOST_DEVICE

903LongIndex capacity(MatrixCoord const &extent) const {

904Index s;

905if (layout_id_ == MatrixLayout::kRowMajor) {

906 s = extent.row();

907 }

908else {

909 s = extent.column();

910 }

911

912Index v = Index((s + stride_[1] - 1) / stride_[1]);

913return LongIndex(v) * LongIndex(stride_[0]);

914 }

915 };

916

918

920 template <typename Layout>

921 struct LayoutTranspose;

922

924 template <>

925 struct LayoutTranspose<layout::RowMajor> {

926using type = layout::ColumnMajor;

927 };

928

930 template <>

931 struct LayoutTranspose<layout::ColumnMajor> {

932using type = layout::RowMajor;

933 };

934

936

937 } // namespace layout

938 } // namespace cutlass

cutlass::layout::ContiguousMatrix::Index

int32_t Index

Index type used for coordinates.

Definition: layout/matrix.h:464

cutlass::layout::RowMajorInterleaved::packed

static CUTLASS_HOST_DEVICE RowMajorInterleaved packed(MatrixCoord const &extent)

Helper returns a layout to a tightly packed tensor.

Definition: layout/matrix.h:283

cutlass::layout::RowMajorInterleaved::stride

CUTLASS_HOST_DEVICE Index stride(int idx) const

Returns the stride of the layout.

Definition: layout/matrix.h:323

cutlass::layout::RowMajor::LongIndex

int64_t LongIndex

Long index type used for offsets.

Definition: layout/matrix.h:62

cutlass::layout::ColumnMajorBlockLinear::stride

CUTLASS_HOST_DEVICE Index & stride(int idx)

Returns the stride of the layout.

Definition: layout/matrix.h:668

cutlass::layout::ColumnMajorInterleaved< 4 >::LongIndex

int64_t LongIndex

Long index type used for offsets.

Definition: layout/matrix.h:355

cutlass::MatrixCoord::column

CUTLASS_HOST_DEVICE Index const & column() const

Returns the column of the coordinate.

Definition: matrix_coord.h:85

cutlass::layout::ContiguousMatrix::stride

CUTLASS_HOST_DEVICE Stride stride() const

Returns the stride of the layout.

Definition: layout/matrix.h:540

cutlass::layout::RowMajorBlockLinear::packed

static CUTLASS_HOST_DEVICE RowMajorBlockLinear packed(MatrixCoord const &extent)

Helper returns a layout to a tightly packed tensor.

Definition: layout/matrix.h:726

cutlass

Definition: aligned_buffer.h:35

cutlass::layout::RowMajorBlockLinear::stride

CUTLASS_HOST_DEVICE Stride stride() const

Returns the stride of the layout.

Definition: layout/matrix.h:750

cutlass::layout::Matrix::kColumnMajor

leading dimension refers to stride between columns; stride along rows is 1

cutlass::layout::RowMajorInterleaved::stride

CUTLASS_HOST_DEVICE Stride stride() const

Returns the stride of the layout.

Definition: layout/matrix.h:311

cutlass::layout::RowMajorBlockLinear::stride

CUTLASS_HOST_DEVICE Index & stride(int idx)

Returns the stride of the layout.

Definition: layout/matrix.h:768

cutlass::layout::ColumnMajorInterleaved::ColumnMajorInterleaved

CUTLASS_HOST_DEVICE ColumnMajorInterleaved(Index ldm=0)

Ctor.

Definition: layout/matrix.h:381

cutlass::layout::GeneralMatrix::GeneralMatrix

CUTLASS_HOST_DEVICE GeneralMatrix(MatrixLayout layout_id, Index ldm, Index interleave)

Ctor.

Definition: layout/matrix.h:822

cutlass::layout::ColumnMajorBlockLinear::LongIndex

int64_t LongIndex

Long index type used for offsets.

Definition: layout/matrix.h:592

cutlass::layout::ContiguousMatrix

Definition: layout/matrix.h:455

cutlass::layout::GeneralMatrix

Definition: layout/matrix.h:781

cutlass::layout::ColumnMajorBlockLinear::stride

CUTLASS_HOST_DEVICE Index stride(int idx) const

Returns the stride of the layout.

Definition: layout/matrix.h:662

cutlass::layout::ColumnMajorInterleaved::capacity

CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const

Compute the number of contiguous elements needed to store a tensor with the given size...

Definition: layout/matrix.h:442

cutlass::layout::RowMajor::stride

CUTLASS_HOST_DEVICE Index & stride(int idx)

Returns the stride of the layout.

Definition: layout/matrix.h:130

cutlass::layout::RowMajor::stride

CUTLASS_HOST_DEVICE Index stride(int idx) const

Returns the stride of the layout.

Definition: layout/matrix.h:124

cutlass::layout::ContiguousMatrix::LongIndex

int64_t LongIndex

Long index type used for offsets.

Definition: layout/matrix.h:467

cutlass::layout::RowMajorInterleaved< 4 >::LongIndex

int64_t LongIndex

Long index type used for offsets.

Definition: layout/matrix.h:249

cutlass::layout::ColumnMajorInterleaved::operator()

CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const

Definition: layout/matrix.h:397

cutlass::layout::GeneralMatrix::LongIndex

int64_t LongIndex

Long index type used for offsets.

Definition: layout/matrix.h:793

cutlass::layout::RowMajorBlockLinear::LongIndex

int64_t LongIndex

Long index type used for offsets.

Definition: layout/matrix.h:693

cutlass::layout::RowMajor::kRank

static int const kRank

Logical rank of tensor.

Definition: layout/matrix.h:53

cutlass::layout::ColumnMajor::stride

CUTLASS_HOST_DEVICE Stride & stride()

Returns the stride of the layout.

Definition: layout/matrix.h:211

cutlass::layout::RowMajorInterleaved::stride

CUTLASS_HOST_DEVICE Index & stride(int idx)

Returns the stride of the layout.

Definition: layout/matrix.h:329

cutlass::layout::ColumnMajorInterleaved< 4 >::Index

int32_t Index

Index type used for coordinates.

Definition: layout/matrix.h:352

cutlass::layout::ContiguousMatrix::stride

CUTLASS_HOST_DEVICE Index & stride(int idx)

Returns the stride of the layout.

Definition: layout/matrix.h:558

cutlass::make_Coord

CUTLASS_HOST_DEVICE Coord< 1 > make_Coord(int _0)

Helper to make a 2-element coordinate.

Definition: coord.h:387

cutlass::layout::ColumnMajor::LongIndex

int64_t LongIndex

Long index type used for offsets.

Definition: layout/matrix.h:154

cutlass::layout::ColumnMajor::ColumnMajor

CUTLASS_HOST_DEVICE ColumnMajor(Index ldm=0)

Ctor.

Definition: layout/matrix.h:177

cutlass::layout::RowMajorInterleaved::inverse

CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const

Inverse of layout function, mapping linear offset to logical coordinate.

Definition: layout/matrix.h:298

cutlass::layout::RowMajorBlockLinear::operator()

CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const

Definition: layout/matrix.h:733

cutlass::layout::RowMajorInterleaved< 4 >::Index

int32_t Index

Index type used for coordinates.

Definition: layout/matrix.h:246

cutlass::MatrixCoord::row

CUTLASS_HOST_DEVICE Index const & row() const

Returns the row of the coordinate.

Definition: matrix_coord.h:77

cutlass::layout::RowMajor::capacity

CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const

Compute the number of contiguous elements needed to store a tensor with the given size...

Definition: layout/matrix.h:136

cutlass::layout::ColumnMajorInterleaved::stride

CUTLASS_HOST_DEVICE Stride stride() const

Returns the stride of the layout.

Definition: layout/matrix.h:418

cutlass::layout::RowMajor::stride

CUTLASS_HOST_DEVICE Stride stride() const

Returns the stride of the layout.

Definition: layout/matrix.h:112

cutlass::layout::RowMajorInterleaved::capacity

CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const

Compute the number of contiguous elements needed to store a tensor with the given size...

Definition: layout/matrix.h:335

cutlass::layout::RowMajor::inverse

CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const

Inverse of layout function, mapping linear offset to logical coordinate.

Definition: layout/matrix.h:106

cutlass::layout::ColumnMajorBlockLinear

Definition: layout/matrix.h:581

cutlass::layout::ColumnMajor

Mapping function for column-major matrices.

Definition: layout/matrix.h:142

cutlass::layout::GeneralMatrix::layout_id

CUTLASS_HOST_DEVICE MatrixLayout & layout_id()

Definition: layout/matrix.h:885

cutlass::layout::RowMajorBlockLinear::stride

CUTLASS_HOST_DEVICE Index stride(int idx) const

Returns the stride of the layout.

Definition: layout/matrix.h:762

cutlass::layout::GeneralMatrix::stride

CUTLASS_HOST_DEVICE Index stride(int idx) const

Returns the stride of the layout.

Definition: layout/matrix.h:891

cutlass::layout::GeneralMatrix::GeneralMatrix

CUTLASS_HOST_DEVICE GeneralMatrix()

Ctor.

Definition: layout/matrix.h:818

cutlass::layout::RowMajor::Index

int32_t Index

Index type used for coordinates.

Definition: layout/matrix.h:59

cutlass::layout::ColumnMajorBlockLinear::stride

CUTLASS_HOST_DEVICE Stride stride() const

Returns the stride of the layout.

Definition: layout/matrix.h:650

cutlass::layout::ColumnMajorBlockLinear::inverse

CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const

Inverse of layout function, mapping linear offset to logical coordinate.

Definition: layout/matrix.h:642

cutlass::layout::GeneralMatrix::stride

CUTLASS_HOST_DEVICE Index & stride(int idx)

Returns the stride of the layout.

Definition: layout/matrix.h:897

cutlass::layout::RowMajor::operator()

CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const

Definition: layout/matrix.h:100

cutlass::layout::GeneralMatrix::Index

int32_t Index

Index type used for coordinates.

Definition: layout/matrix.h:790

cutlass::layout::LayoutTranspose

Defines transposes of matrix layouts.

Definition: layout/matrix.h:921

cutlass::layout::ColumnMajorBlockLinear::operator()

CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const

Definition: layout/matrix.h:632

cutlass::layout::ContiguousMatrix::capacity

CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const

Compute the number of contiguous elements needed to store a tensor with the given size...

Definition: layout/matrix.h:564

cutlass::layout::ColumnMajorBlockLinear::Index

int32_t Index

Index type used for coordinates.

Definition: layout/matrix.h:589

cutlass::layout::Matrix

Matrix

Enumerated type for canonical pitch-linear matrix layouts.

Definition: layout/matrix.h:448

cutlass::layout::ContiguousMatrix::operator()

CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const

Definition: layout/matrix.h:518

cutlass::layout::ColumnMajor::inverse

CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const

Inverse of layout function, mapping linear offset to logical coordinate.

Definition: layout/matrix.h:199

cutlass::layout::RowMajorBlockLinear

Definition: layout/matrix.h:682

cutlass::MatrixLayout

MatrixLayout

Definition: matrix_traits.h:36

cutlass::layout::ColumnMajorInterleaved::stride

CUTLASS_HOST_DEVICE Index & stride(int idx)

Returns the stride of the layout.

Definition: layout/matrix.h:436

cutlass::layout::RowMajorBlockLinear::Index

int32_t Index

Index type used for coordinates.

Definition: layout/matrix.h:690

cutlass::layout::RowMajorInterleaved::RowMajorInterleaved

CUTLASS_HOST_DEVICE RowMajorInterleaved(Stride stride)

Ctor.

Definition: layout/matrix.h:279

CUTLASS_HOST_DEVICE

#define CUTLASS_HOST_DEVICE

Definition: cutlass.h:89

cutlass::layout::ContiguousMatrix::inverse

CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const

Inverse of layout function, mapping linear offset to logical coordinate.

Definition: layout/matrix.h:533

cutlass::layout::RowMajorBlockLinear::inverse

CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const

Inverse of layout function, mapping linear offset to logical coordinate.

Definition: layout/matrix.h:743

cutlass::layout::ColumnMajor::packed

static CUTLASS_HOST_DEVICE ColumnMajor packed(MatrixCoord const &extent)

Helper returns a layout to a tightly packed tensor.

Definition: layout/matrix.h:186

cutlass::layout::RowMajorBlockLinear::stride

CUTLASS_HOST_DEVICE Stride & stride()

Returns the stride of the layout.

Definition: layout/matrix.h:756

cutlass::layout::RowMajorBlockLinear::capacity

CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const

Compute the number of contiguous elements needed to store a tensor with the given size...

Definition: layout/matrix.h:774

cutlass::layout::GeneralMatrix::packed

static CUTLASS_HOST_DEVICE GeneralMatrix packed(MatrixCoord const &extent, MatrixLayout layout_id=MatrixLayout::kColumnMajor, Index interleave=1)

Helper returns a layout to a tightly packed tensor.

Definition: layout/matrix.h:829

cutlass::layout::ColumnMajor::capacity

CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const

Compute the number of contiguous elements needed to store a tensor with the given size...

Definition: layout/matrix.h:229

cutlass::layout::RowMajor::stride

CUTLASS_HOST_DEVICE Stride & stride()

Returns the stride of the layout.

Definition: layout/matrix.h:118

cutlass::layout::ContiguousMatrix::packed

static CUTLASS_HOST_DEVICE ContiguousMatrix packed(MatrixCoord const &extent, Matrix layout=Matrix::kColumnMajor)

Helper returns a layout to a tightly packed tensor.

Definition: layout/matrix.h:501

cutlass::layout::ColumnMajor::stride

CUTLASS_HOST_DEVICE Stride stride() const

Returns the stride of the layout.

Definition: layout/matrix.h:205

cutlass::layout::ColumnMajorInterleaved::stride

CUTLASS_HOST_DEVICE Stride & stride()

Returns the stride of the layout.

Definition: layout/matrix.h:424

cutlass::layout::Matrix::kRowMajor

leading dimension refers to stride between rows; stride along columns is 1

cutlass::Coord< kStrideRank, Index >

cutlass::layout::RowMajor

Mapping function for row-major matrices.

Definition: layout/matrix.h:50

cutlass::layout::ColumnMajor::operator()

CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const

Definition: layout/matrix.h:193

matrix_coord.h

Defines a canonical coordinate for rank=2 matrices offering named indices.

cutlass::layout::RowMajor::RowMajor

CUTLASS_HOST_DEVICE RowMajor(Stride stride)

Ctor.

Definition: layout/matrix.h:89

cutlass::layout::GeneralMatrix::stride

CUTLASS_HOST_DEVICE Stride & stride()

Returns the stride of the layout.

Definition: layout/matrix.h:880

cutlass::layout::RowMajorInterleaved::operator()

CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const

Definition: layout/matrix.h:290

cutlass::layout::ContiguousMatrix::ContiguousMatrix

CUTLASS_HOST_DEVICE ContiguousMatrix(Index ldm=0, Matrix layout=Matrix::kColumnMajor)

Ctor.

Definition: layout/matrix.h:493

cutlass::layout::ColumnMajorInterleaved::stride

CUTLASS_HOST_DEVICE Index stride(int idx) const

Returns the stride of the layout.

Definition: layout/matrix.h:430

cutlass::layout::RowMajorInterleaved::stride

CUTLASS_HOST_DEVICE Stride & stride()

Returns the stride of the layout.

Definition: layout/matrix.h:317

cutlass::layout::ContiguousMatrix::stride

CUTLASS_HOST_DEVICE Index stride(int idx) const

Returns the stride of the layout.

Definition: layout/matrix.h:552

cutlass::layout::ColumnMajor::stride

CUTLASS_HOST_DEVICE Index stride(int idx) const

Returns the stride of the layout.

Definition: layout/matrix.h:217

cutlass::layout::ColumnMajorBlockLinear::packed

static CUTLASS_HOST_DEVICE ColumnMajorBlockLinear packed(MatrixCoord const &extent)

Helper returns a layout to a tightly packed tensor.

Definition: layout/matrix.h:625

cutlass::layout::ColumnMajorInterleaved::ColumnMajorInterleaved

CUTLASS_HOST_DEVICE ColumnMajorInterleaved(Stride stride)

Ctor.

Definition: layout/matrix.h:385

cutlass::layout::GeneralMatrix::operator()

CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const

Definition: layout/matrix.h:850

cutlass::layout::ColumnMajorInterleaved

Definition: layout/matrix.h:343

cutlass::layout::ColumnMajor::Index

int32_t Index

Index type used for coordinates.

Definition: layout/matrix.h:151

cutlass::layout::GeneralMatrix::stride

CUTLASS_HOST_DEVICE Stride stride() const

Returns the stride of the layout.

Definition: layout/matrix.h:869

cutlass::layout::ColumnMajorBlockLinear::stride

CUTLASS_HOST_DEVICE Stride & stride()

Returns the stride of the layout.

Definition: layout/matrix.h:656

cutlass::layout::RowMajor::packed

static CUTLASS_HOST_DEVICE RowMajor packed(MatrixCoord const &extent)

Helper returns a layout to a tightly packed tensor.

Definition: layout/matrix.h:93

cutlass::layout::RowMajorInterleaved::RowMajorInterleaved

CUTLASS_HOST_DEVICE RowMajorInterleaved(Index ldm=0)

Ctor.

Definition: layout/matrix.h:275

cutlass::layout::RowMajorBlockLinear::RowMajorBlockLinear

CUTLASS_HOST_DEVICE RowMajorBlockLinear(Index ldm=0)

Ctor.

Definition: layout/matrix.h:722

cutlass::layout::GeneralMatrix::capacity

CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const

Compute the number of contiguous elements needed to store a tensor with the given size...

Definition: layout/matrix.h:903

cutlass::layout::ColumnMajorInterleaved::inverse

CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const

Inverse of layout function, mapping linear offset to logical coordinate.

Definition: layout/matrix.h:405

matrix_traits.h

Defines properties of matrices used to denote layout and operands to GEMM kernels.

cutlass::layout::ColumnMajorInterleaved::packed

static CUTLASS_HOST_DEVICE ColumnMajorInterleaved packed(MatrixCoord const &extent)

Helper returns a layout to a tightly packed tensor.

Definition: layout/matrix.h:390

cutlass::layout::ColumnMajor::stride

CUTLASS_HOST_DEVICE Index & stride(int idx)

Returns the stride of the layout.

Definition: layout/matrix.h:223

cutlass::layout::ContiguousMatrix::stride

CUTLASS_HOST_DEVICE Stride & stride()

Returns the stride of the layout.

Definition: layout/matrix.h:546

cutlass.h

Basic include for CUTLASS.

cutlass::MatrixCoord

Definition: matrix_coord.h:39

cutlass::layout::ColumnMajorBlockLinear::capacity

CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const

Compute the number of contiguous elements needed to store a tensor with the given size...

Definition: layout/matrix.h:674

cutlass::layout::ColumnMajor::ColumnMajor

CUTLASS_HOST_DEVICE ColumnMajor(Stride stride)

Ctor.

Definition: layout/matrix.h:181

cutlass::layout::RowMajor::kStrideRank

static int const kStrideRank

Rank of stride vector.

Definition: layout/matrix.h:56

cutlass::layout::ColumnMajorBlockLinear::ColumnMajorBlockLinear

CUTLASS_HOST_DEVICE ColumnMajorBlockLinear(Index ldm=0)

Ctor.

Definition: layout/matrix.h:621

cutlass::layout::RowMajor::RowMajor

CUTLASS_HOST_DEVICE RowMajor(Index ldm=0)

Constructor.

Definition: layout/matrix.h:85

cutlass::layout::GeneralMatrix::layout_id

CUTLASS_HOST_DEVICE MatrixLayout layout_id() const

Definition: layout/matrix.h:874

cutlass::layout::RowMajorInterleaved

Definition: layout/matrix.h:237

<!-- fragment --> <!-- contents --><!-- start footer part -->
<address class="footer"><small> Generated by 1.8.11 </small></address>