Back to Cutlass

CUTLASS: subbyte_reference.h Source File

docs/subbyte__reference_8h_source.html

4.4.242.3 KB
Original Source

| | CUTLASS

CUDA Templates for Linear Algebra Subroutines and Solvers |

subbyte_reference.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 **************************************************************************************************/

28 #pragma once

29

30 #include "cutlass/numeric_types.h"

31

32 namespace cutlass {

33

35

55 template <

56typename Element_,

57typename Storage_ = uint8_t

58 >

60 class ConstSubbyteReference {

61 public:

62

63using Element = Element_;

64using Storage = Storage_;

65using StoragePointer = Storage const *;

66

67static_assert(sizeof_bits<Element>::value <= sizeof_bits<Storage>::value,

68"Size of Element must not be greater than Storage.");

69

70static_assert(!(sizeof_bits<Storage>::value % sizeof_bits<Element>::value),

71"Storage must be divisible by Element");

72

73 private:

74

76int const kElementsPerVector = sizeof_bits<Storage>::value / sizeof_bits<Element>::value;

77

79Storage const kMask =

80 ((sizeof_bits<Element>::value < sizeof_bits<Storage>::value) ?

81 (Storage(1) << sizeof_bits<Element>::value) - Storage(1) :

82 ~Storage(0));

83

84 private:

85

87StoragePointer ptr_;

88

92int offset_;

93

94 public:

95

96CUTLASS_HOST_DEVICE

97ConstSubbyteReference(): ptr_(nullptr), offset_(0) { }

98

100CUTLASS_HOST_DEVICE

101ConstSubbyteReference(

102Element const *ptr,

103 int64_t offset

104 ):

105 ptr_(reinterpret_cast<StoragePointer>(ptr)),

106 offset_(0) {

107

108 int64_t offset_in_vectors = offset / kElementsPerVector;

109 int64_t offset_in_elements = offset % kElementsPerVector;

110

111 ptr_ += offset_in_vectors;

112 offset_ = int(offset_in_elements);

113 }

114

116CUTLASS_HOST_DEVICE

117ConstSubbyteReference(

118Element *ptr = nullptr

119 ): ConstSubbyteReference(ptr, 0) { }

120

122CUTLASS_HOST_DEVICE

123StoragePointer storage_pointer() const {

124return ptr_;

125 }

126

128CUTLASS_HOST_DEVICE

129int element_offset() const {

130return offset_;

131 }

132

134CUTLASS_HOST_DEVICE

135Element get() const {

136Storage item = Storage((*ptr_ >> (offset_ * sizeof_bits<Element>::value)) & kMask);

137return reinterpret_cast<Element const &>(item);

138 }

139

141CUTLASS_HOST_DEVICE

142operator Element() const {

143return get();

144 }

145

147CUTLASS_HOST_DEVICE

148ConstSubbyteReference &operator+=(int offset) {

149

150 offset += offset_;

151

152int offset_in_vectors = offset / kElementsPerVector;

153int offset_in_elements = offset % kElementsPerVector;

154

155 ptr_ += offset_in_vectors;

156 offset_ = offset_in_elements;

157

158return *this;

159 }

160

162CUTLASS_HOST_DEVICE

163ConstSubbyteReference &operator+=(long long offset) {

164

165 offset += offset_;

166

167long long offset_in_vectors = offset / kElementsPerVector;

168int offset_in_elements = int(offset % kElementsPerVector);

169

170 ptr_ += offset_in_vectors;

171 offset_ = offset_in_elements;

172

173return *this;

174 }

175

177CUTLASS_HOST_DEVICE

178ConstSubbyteReference &operator-=(int offset) {

179

180int offset_in_vectors = offset / kElementsPerVector;

181int offset_in_elements = offset % kElementsPerVector;

182

183 ptr_ -= offset_in_vectors;

184 offset_ -= offset_in_elements;

185

186if (offset_ < 0) {

187 offset_ += kElementsPerVector;

188 --ptr_;

189 }

190

191return *this;

192 }

193

195CUTLASS_HOST_DEVICE

196ConstSubbyteReference &operator-=(long long offset) {

197

198long long offset_in_vectors = offset / kElementsPerVector;

199int offset_in_elements = int(offset % kElementsPerVector);

200

201 ptr_ -= offset_in_vectors;

202 offset_ -= offset_in_elements;

203

204if (offset_ < 0) {

205 offset_ += kElementsPerVector;

206 --ptr_;

207 }

208

209return *this;

210 }

211

213CUTLASS_HOST_DEVICE

214ConstSubbyteReference operator+(int offset) const {

215

216ConstSubbyteReference ref(ptr_, offset_);

217 ref += offset;

218

219return ref;

220 }

221

223CUTLASS_HOST_DEVICE

224ConstSubbyteReference operator+(long long offset) const {

225

226ConstSubbyteReference ref(ptr_, offset_);

227 ref += offset;

228

229return ref;

230 }

231

233CUTLASS_HOST_DEVICE

234ConstSubbyteReference operator-(int offset) const {

235

236ConstSubbyteReference ref(ptr_, offset_);

237 ref -= offset;

238

239return ref;

240 }

241

243CUTLASS_HOST_DEVICE

244ConstSubbyteReference operator-=(long long offset) const {

245

246ConstSubbyteReference ref(ptr_, offset_);

247 ref -= offset;

248

249return ref;

250 }

251

253CUTLASS_HOST_DEVICE

254 ptrdiff_t operator-(ConstSubbyteReference ref) const {

255return (ptr_ - ref.ptr_) * kElementsPerVector + (offset_ - ref.offset_);

256 }

257

259CUTLASS_HOST_DEVICE

260explicit operator int() const {

261return int(get());

262 }

263

265CUTLASS_HOST_DEVICE

266explicit operator int64_t() const {

267return int64_t(get());

268 }

269

271CUTLASS_HOST_DEVICE

272explicit operator uint64_t() const {

273return uint64_t(get());

274 }

275

277CUTLASS_HOST_DEVICE

278explicit operator float() const {

279return float(get());

280 }

281

283CUTLASS_HOST_DEVICE

284explicit operator double() const {

285return double(get());

286 }

287 };

288

289 template <

290typename Element_,

291typename Storage_ = uint8_t

292 >

294 class SubbyteReference {

295 public:

296

297using Element = Element_;

298using Storage = Storage_;

299using StoragePointer = Storage *;

300

301static_assert(sizeof_bits<Element>::value <= sizeof_bits<Storage>::value,

302"Size of Element must not be greater than Storage.");

303

304static_assert(!(sizeof_bits<Storage>::value % sizeof_bits<Element>::value),

305"Storage must be divisible by Element");

306

307 private:

308

310int const kElementsPerVector = sizeof_bits<Storage>::value / sizeof_bits<Element>::value;

311

313Storage const kMask =

314 ((sizeof_bits<Element>::value < sizeof_bits<Storage>::value) ?

315 (Storage(1) << sizeof_bits<Element>::value) - Storage(1) :

316 ~Storage(0));

317

318 private:

319

321StoragePointer ptr_;

322

326int offset_;

327

328 public:

329

330CUTLASS_HOST_DEVICE

331SubbyteReference(): ptr_(nullptr), offset_(0) { }

332

334CUTLASS_HOST_DEVICE

335SubbyteReference(

336Element *ptr,

337 int64_t offset

338 ):

339 ptr_(reinterpret_cast<StoragePointer>(ptr)),

340 offset_(0) {

341

342 int64_t offset_in_vectors = offset / kElementsPerVector;

343 int64_t offset_in_elements = offset % kElementsPerVector;

344

345 ptr_ += offset_in_vectors;

346 offset_ = int(offset_in_elements);

347 }

348

350CUTLASS_HOST_DEVICE

351SubbyteReference(

352Element *ptr = nullptr

353 ): SubbyteReference(ptr, 0) { }

354

356CUTLASS_HOST_DEVICE

357StoragePointer storage_pointer() const {

358return ptr_;

359 }

360

362CUTLASS_HOST_DEVICE

363int element_offset() const {

364return offset_;

365 }

366

368CUTLASS_HOST_DEVICE

369Element get() const {

370Storage item = Storage((*ptr_ >> (offset_ * sizeof_bits<Element>::value)) & kMask);

371return reinterpret_cast<Element const &>(item);

372 }

373

375CUTLASS_HOST_DEVICE

376SubbyteReference & set(Element const &x) {

377

378Storage item = (reinterpret_cast<Storage const &>(x) & kMask);

379

380Storage kUpdateMask = Storage(~(kMask << (offset_ * sizeof_bits<Element>::value)));

381 *ptr_ = Storage((*ptr_ & kUpdateMask) | Storage(item << (offset_ * sizeof_bits<Element>::value)));

382

383return *this;

384 }

385

387CUTLASS_HOST_DEVICE

388operator Element() const {

389return get();

390 }

391

393CUTLASS_HOST_DEVICE

394SubbyteReference &operator=(Element const & x) {

395return set(x);

396 }

397

399CUTLASS_HOST_DEVICE

400SubbyteReference &operator=(SubbyteReference const & x) {

401return set(x.get());

402 }

403

405CUTLASS_HOST_DEVICE

406SubbyteReference &operator=(

407ConstSubbyteReference<Element, Storage> const &x) {

408return set(x.get());

409 }

410

412CUTLASS_HOST_DEVICE

413SubbyteReference &operator+=(int offset) {

414

415 offset += offset_;

416

417int offset_in_vectors = offset / kElementsPerVector;

418int offset_in_elements = offset % kElementsPerVector;

419

420 ptr_ += offset_in_vectors;

421 offset_ = offset_in_elements;

422

423return *this;

424 }

425

427CUTLASS_HOST_DEVICE

428SubbyteReference &operator+=(long long offset) {

429

430 offset += offset_;

431

432long long offset_in_vectors = offset / kElementsPerVector;

433int offset_in_elements = int(offset % kElementsPerVector);

434

435 ptr_ += offset_in_vectors;

436 offset_ = offset_in_elements;

437

438return *this;

439 }

440

442CUTLASS_HOST_DEVICE

443SubbyteReference &operator-=(int offset) {

444

445int offset_in_vectors = offset / kElementsPerVector;

446int offset_in_elements = offset % kElementsPerVector;

447

448 ptr_ -= offset_in_vectors;

449 offset_ -= offset_in_elements;

450

451if (offset_ < 0) {

452 offset_ += kElementsPerVector;

453 --ptr_;

454 }

455

456return *this;

457 }

458

460CUTLASS_HOST_DEVICE

461SubbyteReference &operator-=(long long offset) {

462

463long long offset_in_vectors = offset / kElementsPerVector;

464int offset_in_elements = int(offset % kElementsPerVector);

465

466 ptr_ -= offset_in_vectors;

467 offset_ -= offset_in_elements;

468

469if (offset_ < 0) {

470 offset_ += kElementsPerVector;

471 --ptr_;

472 }

473

474return *this;

475 }

476

478CUTLASS_HOST_DEVICE

479SubbyteReference operator+(int offset) const {

480

481SubbyteReference ref(ptr_, offset_);

482 ref += offset;

483

484return ref;

485 }

486

488CUTLASS_HOST_DEVICE

489SubbyteReference operator+(long long offset) const {

490

491SubbyteReference ref(ptr_, offset_);

492 ref += offset;

493

494return ref;

495 }

496

498CUTLASS_HOST_DEVICE

499SubbyteReference operator-(int offset) const {

500

501SubbyteReference ref(ptr_, offset_);

502 ref -= offset;

503

504return ref;

505 }

506

508CUTLASS_HOST_DEVICE

509SubbyteReference operator-=(long long offset) const {

510

511SubbyteReference ref(ptr_, offset_);

512 ref -= offset;

513

514return ref;

515 }

516

518CUTLASS_HOST_DEVICE

519 ptrdiff_t operator-(SubbyteReference ref) const {

520return (ptr_ - ref.ptr_) * kElementsPerVector + (offset_ - ref.offset_);

521 }

522

524CUTLASS_HOST_DEVICE

525explicit operator int() const {

526return int(get());

527 }

528

530CUTLASS_HOST_DEVICE

531explicit operator int64_t() const {

532return int64_t(get());

533 }

534

536CUTLASS_HOST_DEVICE

537explicit operator uint64_t() const {

538return uint64_t(get());

539 }

540

542CUTLASS_HOST_DEVICE

543explicit operator float() const {

544return float(get());

545 }

546

548CUTLASS_HOST_DEVICE

549explicit operator double() const {

550return double(get());

551 }

552 };

553

555

556 template <typename Element, bool subbyte = (sizeof_bits<Element>::value < 8)>

557 struct ReferenceFactory;

558

559 template <typename Element>

560 struct ReferenceFactory<Element, false> {

561CUTLASS_HOST_DEVICE

562static Element &get(Element *ptr, int64_t offset) {

563return ptr[offset];

564 }

565

566CUTLASS_HOST_DEVICE

567static Element const &get(Element const *ptr, int64_t offset) {

568return ptr[offset];

569 }

570 };

571

572 template <typename Element>

573 struct ReferenceFactory<Element, true> {

574CUTLASS_HOST_DEVICE

575static SubbyteReference<Element> get(Element *ptr, int64_t offset) {

576return SubbyteReference<Element>(ptr, offset);

577 }

578

579CUTLASS_HOST_DEVICE

580static ConstSubbyteReference<Element> get(Element const *ptr,

581 int64_t offset) {

582return ConstSubbyteReference<Element>(ptr, offset);

583 }

584 };

585

587

588 } // namespace cutlass

cutlass::ConstSubbyteReference

Definition: subbyte_reference.h:60

cutlass

Definition: aligned_buffer.h:35

cutlass::SubbyteReference::element_offset

CUTLASS_HOST_DEVICE int element_offset() const

Gets element offset within storage vector.

Definition: subbyte_reference.h:363

cutlass::SubbyteReference::operator=

CUTLASS_HOST_DEVICE SubbyteReference & operator=(Element const &x)

Stores an element to memory.

Definition: subbyte_reference.h:394

cutlass::SubbyteReference::SubbyteReference

CUTLASS_HOST_DEVICE SubbyteReference()

Definition: subbyte_reference.h:331

cutlass::SubbyteReference::operator-=

CUTLASS_HOST_DEVICE SubbyteReference operator-=(long long offset) const

Returns a reference to an element with a given offset from the current reference. ...

Definition: subbyte_reference.h:509

cutlass::SubbyteReference::SubbyteReference

CUTLASS_HOST_DEVICE SubbyteReference(Element *ptr, int64_t offset)

Constructor.

Definition: subbyte_reference.h:335

cutlass::SubbyteReference::operator+

CUTLASS_HOST_DEVICE SubbyteReference operator+(int offset) const

Returns a reference to an element with a given offset from the current reference. ...

Definition: subbyte_reference.h:479

cutlass::ConstSubbyteReference::operator-=

CUTLASS_HOST_DEVICE ConstSubbyteReference operator-=(long long offset) const

Returns a reference to an element with a given offset from the current reference. ...

Definition: subbyte_reference.h:244

cutlass::ConstSubbyteReference::Storage

Storage_ Storage

Definition: subbyte_reference.h:64

cutlass::ConstSubbyteReference::operator+=

CUTLASS_HOST_DEVICE ConstSubbyteReference & operator+=(long long offset)

Adds an offset in units of elements to the reference.

Definition: subbyte_reference.h:163

cutlass::SubbyteReference::operator+

CUTLASS_HOST_DEVICE SubbyteReference operator+(long long offset) const

Returns a reference to an element with a given offset from the current reference. ...

Definition: subbyte_reference.h:489

cutlass::ConstSubbyteReference::operator-

CUTLASS_HOST_DEVICE ptrdiff_t operator-(ConstSubbyteReference ref) const

Computes the difference in elements between references.

Definition: subbyte_reference.h:254

cutlass::ConstSubbyteReference::operator+=

CUTLASS_HOST_DEVICE ConstSubbyteReference & operator+=(int offset)

Adds an offset in units of elements to the reference.

Definition: subbyte_reference.h:148

cutlass::SubbyteReference::operator-=

CUTLASS_HOST_DEVICE SubbyteReference & operator-=(long long offset)

Adds an offset in units of elements to the reference.

Definition: subbyte_reference.h:461

cutlass::sizeof_bits

Defines the size of an element in bits.

Definition: numeric_types.h:42

cutlass::ConstSubbyteReference::get

CUTLASS_HOST_DEVICE Element get() const

Unpacks an element from memory.

Definition: subbyte_reference.h:135

nullptr

#define nullptr

nullptr

Definition: platform.h:144

cutlass::SubbyteReference::Element

Element_ Element

Definition: subbyte_reference.h:297

cutlass::ReferenceFactory

Definition: subbyte_reference.h:557

cutlass::ConstSubbyteReference::storage_pointer

CUTLASS_HOST_DEVICE StoragePointer storage_pointer() const

Gets storage pointer.

Definition: subbyte_reference.h:123

cutlass::SubbyteReference::operator+=

CUTLASS_HOST_DEVICE SubbyteReference & operator+=(long long offset)

Adds an offset in units of elements to the reference.

Definition: subbyte_reference.h:428

CUTLASS_HOST_DEVICE

#define CUTLASS_HOST_DEVICE

Definition: cutlass.h:89

numeric_types.h

Top-level include for all CUTLASS numeric types.

cutlass::SubbyteReference::storage_pointer

CUTLASS_HOST_DEVICE StoragePointer storage_pointer() const

Gets storage pointer.

Definition: subbyte_reference.h:357

static_assert

#define static_assert(__e, __m)

Definition: platform.h:153

cutlass::ConstSubbyteReference::ConstSubbyteReference

CUTLASS_HOST_DEVICE ConstSubbyteReference(Element const *ptr, int64_t offset)

Constructor.

Definition: subbyte_reference.h:101

cutlass::ConstSubbyteReference::operator+

CUTLASS_HOST_DEVICE ConstSubbyteReference operator+(int offset) const

Returns a reference to an element with a given offset from the current reference. ...

Definition: subbyte_reference.h:214

cutlass::ConstSubbyteReference::operator-=

CUTLASS_HOST_DEVICE ConstSubbyteReference & operator-=(int offset)

Adds an offset in units of elements to the reference.

Definition: subbyte_reference.h:178

cutlass::ConstSubbyteReference::ConstSubbyteReference

CUTLASS_HOST_DEVICE ConstSubbyteReference()

Definition: subbyte_reference.h:97

cutlass::ConstSubbyteReference::operator+

CUTLASS_HOST_DEVICE ConstSubbyteReference operator+(long long offset) const

Returns a reference to an element with a given offset from the current reference. ...

Definition: subbyte_reference.h:224

cutlass::SubbyteReference::operator=

CUTLASS_HOST_DEVICE SubbyteReference & operator=(SubbyteReference const &x)

Stores an element to memory.

Definition: subbyte_reference.h:400

cutlass::SubbyteReference::Storage

Storage_ Storage

Definition: subbyte_reference.h:298

cutlass::ConstSubbyteReference::ConstSubbyteReference

CUTLASS_HOST_DEVICE ConstSubbyteReference(Element *ptr=nullptr)

Constructor.

Definition: subbyte_reference.h:117

cutlass::SubbyteReference

Definition: subbyte_reference.h:294

cutlass::ConstSubbyteReference::element_offset

CUTLASS_HOST_DEVICE int element_offset() const

Gets element offset within storage vector.

Definition: subbyte_reference.h:129

cutlass::SubbyteReference::get

CUTLASS_HOST_DEVICE Element get() const

Unpacks an element from memory.

Definition: subbyte_reference.h:369

cutlass::SubbyteReference::operator=

CUTLASS_HOST_DEVICE SubbyteReference & operator=(ConstSubbyteReference< Element, Storage > const &x)

Stores an element to memory.

Definition: subbyte_reference.h:406

cutlass::SubbyteReference::StoragePointer

Storage * StoragePointer

Definition: subbyte_reference.h:299

cutlass::ConstSubbyteReference::StoragePointer

Storage const * StoragePointer

Definition: subbyte_reference.h:65

cutlass::SubbyteReference::SubbyteReference

CUTLASS_HOST_DEVICE SubbyteReference(Element *ptr=nullptr)

Constructor.

Definition: subbyte_reference.h:351

cutlass::SubbyteReference::operator-=

CUTLASS_HOST_DEVICE SubbyteReference & operator-=(int offset)

Adds an offset in units of elements to the reference.

Definition: subbyte_reference.h:443

cutlass::ConstSubbyteReference::Element

Element_ Element

Definition: subbyte_reference.h:63

cutlass::ConstSubbyteReference::operator-=

CUTLASS_HOST_DEVICE ConstSubbyteReference & operator-=(long long offset)

Adds an offset in units of elements to the reference.

Definition: subbyte_reference.h:196

cutlass::SubbyteReference::operator-

CUTLASS_HOST_DEVICE SubbyteReference operator-(int offset) const

Returns a reference to an element with a given offset from the current reference. ...

Definition: subbyte_reference.h:499

cutlass::ConstSubbyteReference::operator-

CUTLASS_HOST_DEVICE ConstSubbyteReference operator-(int offset) const

Returns a reference to an element with a given offset from the current reference. ...

Definition: subbyte_reference.h:234

cutlass::SubbyteReference::operator-

CUTLASS_HOST_DEVICE ptrdiff_t operator-(SubbyteReference ref) const

Computes the difference in elements between references.

Definition: subbyte_reference.h:519

cutlass::SubbyteReference::operator+=

CUTLASS_HOST_DEVICE SubbyteReference & operator+=(int offset)

Adds an offset in units of elements to the reference.

Definition: subbyte_reference.h:413


Generated by 1.8.11