CUTLASS: array_subbyte.h Source File - Cutlass

| | CUTLASS

CUDA Templates for Linear Algebra Subroutines and Solvers |

array_subbyte.h

Go to the documentation of this file.

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

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

30 #pragma once

32 #include "cutlass/cutlass.h"

33 #include "cutlass/array.h"

34 #include "cutlass/platform/platform.h"

36 namespace cutlass {

41 template <

42typename T,

43int N

44 >

45 class Array<T, N, false> {

46 public:

48static_assert(sizeof_bits<T>::value * N >= 8,

49"Array<> specialized for sub-byte types assume the actual stored element size is 1 byte");

51static int const kSizeBits = sizeof_bits<T>::value * N;

54using Storage = typename platform::conditional<

55 ((kSizeBits % 32) != 0),

56typename platform::conditional<

57 ((kSizeBits % 16) != 0),

58 uint8_t,

59 uint16_t

60 >::type,

61 uint32_t

62 >::type;

65using Element = T;

68static int const kElementsPerStoredItem = (sizeof(Storage) * 8) / sizeof_bits<T>::value;

71static size_t const kStorageElements = N / kElementsPerStoredItem;

74static size_t const kElements = N;

77static Storage const kMask = ((Storage(1) << sizeof_bits<T>::value) - 1);

79//

80// C++ standard members with pointer types removed

81//

83typedef T value_type;

84typedef size_t size_type;

85typedef ptrdiff_t difference_type;

86typedef value_type *pointer;

87typedef value_type const *const_pointer;

89//

90// References

91//

94class reference {

96Storage *ptr_;

99int idx_;

100

101public:

102

104CUTLASS_HOST_DEVICE

105 reference(): ptr_(nullptr), idx_(0) { }

106

108CUTLASS_HOST_DEVICE

109 reference(Storage *ptr, int idx = 0): ptr_(ptr), idx_(idx) { }

110

112CUTLASS_HOST_DEVICE

113 reference &operator=(T x) {

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

115

116Storage kUpdateMask = Storage(~(kMask << (idx_ * sizeof_bits<T>::value)));

117 *ptr_ = Storage(((*ptr_ & kUpdateMask) | (item << idx_ * sizeof_bits<T>::value)));

118

119return *this;

120 }

121

122CUTLASS_HOST_DEVICE

123 T get() const {

124Storage item = Storage((*ptr_ >> (idx_ * sizeof_bits<T>::value)) & kMask);

125return reinterpret_cast<T const &>(item);

126 }

127

129CUTLASS_HOST_DEVICE

130operator T() const {

131return get();

132 }

133

135CUTLASS_HOST_DEVICE

136explicit operator int() const {

137return int(get());

138 }

139

141CUTLASS_HOST_DEVICE

142explicit operator float() const {

143return float(get());

144 }

145 };

146

148class const_reference {

149

151Storage const *ptr_;

152

154int idx_;

155

156public:

157

159CUTLASS_HOST_DEVICE

160 const_reference(): ptr_(nullptr), idx_(0) { }

161

163CUTLASS_HOST_DEVICE

164 const_reference(Storage const *ptr, int idx = 0): ptr_(ptr), idx_(idx) { }

165

166CUTLASS_HOST_DEVICE

167const T get() const {

168Storage item = (*ptr_ >> (idx_ * sizeof_bits<T>::value)) & kMask;

169return reinterpret_cast<T const &>(item);

170 }

171

173CUTLASS_HOST_DEVICE

174operator T() const {

175Storage item = Storage(Storage(*ptr_ >> Storage(idx_ * sizeof_bits<T>::value)) & kMask);

176return reinterpret_cast<T const &>(item);

177 }

178

180CUTLASS_HOST_DEVICE

181explicit operator int() const {

182return int(get());

183 }

184

186CUTLASS_HOST_DEVICE

187explicit operator float() const {

188return float(get());

189 }

190 };

191

192//

193// Iterators

194//

195

197class iterator {

198

200Storage *ptr_;

201

203int idx_;

204

205public:

206

207CUTLASS_HOST_DEVICE

208 iterator(): ptr_(nullptr), idx_(0) { }

209

210CUTLASS_HOST_DEVICE

211 iterator(Storage *ptr, int idx = 0): ptr_(ptr), idx_(idx) { }

212

213CUTLASS_HOST_DEVICE

214 iterator &operator++() {

215 ++idx_;

216if (idx_ == kElementsPerStoredItem) {

217 ++ptr_;

218 idx_ = 0;

219 }

220return *this;

221 }

222

223CUTLASS_HOST_DEVICE

224 iterator &operator--() {

225if (!idx_) {

226 --ptr_;

227 idx_ = kElementsPerStoredItem - 1;

228 }

229else {

230 --idx_;

231 }

232return *this;

233 }

234

235CUTLASS_HOST_DEVICE

236 iterator operator++(int) {

237 iterator ret(*this);

238 ++idx_;

239if (idx_ == kElementsPerStoredItem) {

240 ++ptr_;

241 idx_ = 0;

242 }

243return ret;

244 }

245

246CUTLASS_HOST_DEVICE

247 iterator operator--(int) {

248 iterator ret(*this);

249if (!idx_) {

250 --ptr_;

251 idx_ = kElementsPerStoredItem - 1;

252 }

253else {

254 --idx_;

255 }

256return ret;

257 }

258

259CUTLASS_HOST_DEVICE

260 reference operator*() const {

261return reference(ptr_, idx_);

262 }

263

264CUTLASS_HOST_DEVICE

265bool operator==(iterator const &other) const {

266return ptr_ == other.ptr_ && idx_ == other.idx_;

267 }

268

269CUTLASS_HOST_DEVICE

270bool operator!=(iterator const &other) const {

271return !(*this == other);

272 }

273 };

274

276class const_iterator {

277

279Storage const *ptr_;

280

282int idx_;

283

284public:

285

286CUTLASS_HOST_DEVICE

287 const_iterator(): ptr_(nullptr), idx_(0) { }

288

289CUTLASS_HOST_DEVICE

290 const_iterator(Storage const *ptr, int idx = 0): ptr_(ptr), idx_(idx) { }

291

292CUTLASS_HOST_DEVICE

293 iterator &operator++() {

294 ++idx_;

295if (idx_ == kElementsPerStoredItem) {

296 ++ptr_;

297 idx_ = 0;

298 }

299return *this;

300 }

301

302CUTLASS_HOST_DEVICE

303 iterator &operator--() {

304if (!idx_) {

305 --ptr_;

306 idx_ = kElementsPerStoredItem - 1;

307 }

308else {

309 --idx_;

310 }

311return *this;

312 }

313

314CUTLASS_HOST_DEVICE

315 iterator operator++(int) {

316 iterator ret(*this);

317 ++idx_;

318if (idx_ == kElementsPerStoredItem) {

319 ++ptr_;

320 idx_ = 0;

321 }

322return ret;

323 }

324

325CUTLASS_HOST_DEVICE

326 iterator operator--(int) {

327 iterator ret(*this);

328if (!idx_) {

329 --ptr_;

330 idx_ = kElementsPerStoredItem - 1;

331 }

332else {

333 --idx_;

334 }

335return ret;

336 }

337

338CUTLASS_HOST_DEVICE

339 const_reference operator*() const {

340return const_reference(ptr_, idx_);

341 }

342

343CUTLASS_HOST_DEVICE

344bool operator==(iterator const &other) const {

345return ptr_ == other.ptr_ && idx_ == other.idx_;

346 }

347

348CUTLASS_HOST_DEVICE

349bool operator!=(iterator const &other) const {

350return !(*this == other);

351 }

352 };

353

355class reverse_iterator {

356

358Storage *ptr_;

359

361int idx_;

362

363public:

364

365CUTLASS_HOST_DEVICE

366 reverse_iterator(): ptr_(nullptr), idx_(0) { }

367

368CUTLASS_HOST_DEVICE

369 reverse_iterator(Storage *ptr, int idx = 0): ptr_(ptr), idx_(idx) { }

370

371// TODO

372 };

373

[375](classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const reverse iterator.html)class const_reverse_iterator {

376

378Storage const *ptr_;

379

381int idx_;

382

383public:

384

385CUTLASS_HOST_DEVICE

[386](classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const reverse iterator.html#aae7705a26ea52ebd18d5f5809d816ee2)[const_reverse_iterator](classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const reverse iterator.html#aae7705a26ea52ebd18d5f5809d816ee2)(): ptr_(nullptr), idx_(0) { }

387

388CUTLASS_HOST_DEVICE

[389](classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const reverse iterator.html#a4bef88847b70f6bca81dd46bd883373b)[const_reverse_iterator](classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const reverse iterator.html#a4bef88847b70f6bca81dd46bd883373b)(Storage const *ptr, int idx = 0): ptr_(ptr), idx_(idx) { }

390

391// TODO

392 };

393

394 private:

395

397Storage storage[kStorageElements];

398

399 public:

400

401CUTLASS_HOST_DEVICE

402 Array() { }

403

404CUTLASS_HOST_DEVICE

405 Array(Array const &x) {

406CUTLASS_PRAGMA_UNROLL

407for (int i = 0; i < int(kStorageElements); ++i) {

408 storage[i] = x.storage[i];

409 }

410 }

411

413CUTLASS_HOST_DEVICE

414void clear() {

415

416CUTLASS_PRAGMA_UNROLL

417for (int i = 0; i < int(kStorageElements); ++i) {

418 storage[i] = Storage(0);

419 }

420 }

421

422CUTLASS_HOST_DEVICE

423 reference at(size_type pos) {

424return reference(storage + pos / kElementsPerStoredItem, pos % kElementsPerStoredItem);

425 }

426

427CUTLASS_HOST_DEVICE

428 const_reference at(size_type pos) const {

429return const_reference(storage + pos / kElementsPerStoredItem, pos % kElementsPerStoredItem);

430 }

431

432CUTLASS_HOST_DEVICE

433 reference operator[](size_type pos) {

434return at(pos);

435 }

436

437CUTLASS_HOST_DEVICE

438 const_reference operator[](size_type pos) const {

439return at(pos);

440 }

441

442CUTLASS_HOST_DEVICE

443 reference front() {

444return at(0);

445 }

446

447CUTLASS_HOST_DEVICE

448 const_reference front() const {

449return at(0);

450 }

451

452CUTLASS_HOST_DEVICE

453 reference back() {

454return reference(storage + kStorageElements - 1, kElementsPerStoredItem - 1);

455 }

456

457CUTLASS_HOST_DEVICE

458 const_reference back() const {

459return const_reference(storage + kStorageElements - 1, kElementsPerStoredItem - 1);

460 }

461

462CUTLASS_HOST_DEVICE

463 pointer data() {

464return reinterpret_cast<pointer>(storage);

465 }

466

467CUTLASS_HOST_DEVICE

468 const_pointer data() const {

469return reinterpret_cast<const_pointer>(storage);

470 }

471

472CUTLASS_HOST_DEVICE

473 Storage * raw_data() {

474return storage;

475 }

476

477CUTLASS_HOST_DEVICE

478 Storage const * raw_data() const {

479return storage;

480 }

481

482

483CUTLASS_HOST_DEVICE

484 constexpr bool empty() const {

485return !kElements;

486 }

487

488CUTLASS_HOST_DEVICE

489 constexpr size_type size() const {

490return kElements;

491 }

492

493CUTLASS_HOST_DEVICE

494 constexpr size_type max_size() const {

495return kElements;

496 }

497

498CUTLASS_HOST_DEVICE

499void fill(T const &value) {

500// TODO

501 }

502

503CUTLASS_HOST_DEVICE

504 iterator begin() {

505return iterator(storage);

506 }

507

508CUTLASS_HOST_DEVICE

509 const_iterator cbegin() const {

510return const_iterator(storage);

511 }

512

513CUTLASS_HOST_DEVICE

514 iterator end() {

515return iterator(storage + kStorageElements);

516 }

517

518CUTLASS_HOST_DEVICE

519 const_iterator cend() const {

520return const_iterator(storage + kStorageElements);

521 }

522

523CUTLASS_HOST_DEVICE

524 reverse_iterator rbegin() {

525return reverse_iterator(storage + kStorageElements);

526 }

527

528CUTLASS_HOST_DEVICE

529 const_reverse_iterator crbegin() const {

530return const_reverse_iterator(storage + kStorageElements);

531 }

532

533CUTLASS_HOST_DEVICE

534 reverse_iterator rend() {

535return reverse_iterator(storage);

536 }

537

538CUTLASS_HOST_DEVICE

539 const_reverse_iterator crend() const {

540return const_reverse_iterator(storage);

541 }

542

543//

544// Comparison operators

545//

546

547 };

548

550

551 } // namespace cutlass

552

cutlass::Array< T, N, false >::const_reference::const_reference

CUTLASS_HOST_DEVICE const_reference(Storage const *ptr, int idx=0)

Ctor.

Definition: array_subbyte.h:164

cutlass::Array< T, N, false >::at

CUTLASS_HOST_DEVICE const_reference at(size_type pos) const

Definition: array_subbyte.h:428

cutlass::Array< T, N, false >::back

CUTLASS_HOST_DEVICE const_reference back() const

Definition: array_subbyte.h:458

cutlass

Definition: aligned_buffer.h:35

cutlass::Array< T, N, false >::reverse_iterator::reverse_iterator

CUTLASS_HOST_DEVICE reverse_iterator()

Definition: array_subbyte.h:366

constexpr

#define constexpr

Definition: platform.h:137

cutlass::Array< T, N, false >::const_iterator::operator--

CUTLASS_HOST_DEVICE iterator operator--(int)

Definition: array_subbyte.h:326

cutlass::Array< T, N, false >::Array

CUTLASS_HOST_DEVICE Array(Array const &x)

Definition: array_subbyte.h:405

cutlass::Array< T, N, false >::operator[]

CUTLASS_HOST_DEVICE reference operator[](size_type pos)

Definition: array_subbyte.h:433

cutlass::Array< T, N, false >::crend

CUTLASS_HOST_DEVICE const_reverse_iterator crend() const

Definition: array_subbyte.h:539

cutlass::Array< T, N, false >::rbegin

CUTLASS_HOST_DEVICE reverse_iterator rbegin()

Definition: array_subbyte.h:524

cutlass::Array< T, N, false >::iterator::operator==

CUTLASS_HOST_DEVICE bool operator==(iterator const &other) const

Definition: array_subbyte.h:265

cutlass::Array< T, N, false >::size_type

size_t size_type

Definition: array_subbyte.h:84

cutlass::Array< T, N, false >::iterator::operator++

CUTLASS_HOST_DEVICE iterator & operator++()

Definition: array_subbyte.h:214

cutlass::Array< T, N, false >::crbegin

CUTLASS_HOST_DEVICE const_reverse_iterator crbegin() const

Definition: array_subbyte.h:529

cutlass::Array< T, N, false >::rend

CUTLASS_HOST_DEVICE reverse_iterator rend()

Definition: array_subbyte.h:534

cutlass::Array< T, N, false >::empty

CUTLASS_HOST_DEVICE constexpr bool empty() const

Definition: array_subbyte.h:484

cutlass::Array< T, N, false >::const_iterator::const_iterator

CUTLASS_HOST_DEVICE const_iterator(Storage const *ptr, int idx=0)

Definition: array_subbyte.h:290

cutlass::Array< T, N, false >::difference_type

ptrdiff_t difference_type

Definition: array_subbyte.h:85

cutlass::Array< T, N, false >::data

CUTLASS_HOST_DEVICE const_pointer data() const

Definition: array_subbyte.h:468

cutlass::Array< T, N, false >::iterator::operator*

CUTLASS_HOST_DEVICE reference operator*() const

Definition: array_subbyte.h:260

platform.h

C++ features that may be otherwise unimplemented for CUDA device functions.

cutlass::Array< T, N, false >::reference::reference

CUTLASS_HOST_DEVICE reference()

Default ctor.

Definition: array_subbyte.h:105

cutlass::Array< T, N, false >::cend

CUTLASS_HOST_DEVICE const_iterator cend() const

Definition: array_subbyte.h:519

cutlass::Array< T, N, false >::reference::reference

CUTLASS_HOST_DEVICE reference(Storage *ptr, int idx=0)

Ctor.

Definition: array_subbyte.h:109

array.h

Statically sized array of elements that accommodates all CUTLASS-supported numeric types and is safe ...

CUTLASS_PRAGMA_UNROLL

#define CUTLASS_PRAGMA_UNROLL

Definition: cutlass.h:110

cutlass::Array< T, N, false >::raw_data

CUTLASS_HOST_DEVICE Storage * raw_data()

Definition: array_subbyte.h:473

cutlass::Array< T, N, false >::const_reference::const_reference

CUTLASS_HOST_DEVICE const_reference()

Default ctor.

Definition: array_subbyte.h:160

cutlass::Array< T, N, false >::back

CUTLASS_HOST_DEVICE reference back()

Definition: array_subbyte.h:453

cutlass::Array< T, N, false >::begin

CUTLASS_HOST_DEVICE iterator begin()

Definition: array_subbyte.h:504

cutlass::sizeof_bits

Defines the size of an element in bits.

Definition: numeric_types.h:42

cutlass::Array< T, N, false >::fill

CUTLASS_HOST_DEVICE void fill(T const &value)

Definition: array_subbyte.h:499

nullptr

#define nullptr

nullptr

Definition: platform.h:144

cutlass::Array< T, N, false >::const_iterator::operator--

CUTLASS_HOST_DEVICE iterator & operator--()

Definition: array_subbyte.h:303

cutlass::Array< T, N, false >::Array

CUTLASS_HOST_DEVICE Array()

Definition: array_subbyte.h:402

cutlass::Array< T, N, false >::Element

T Element

Element type.

Definition: array_subbyte.h:65

cutlass::Array< T, N, false >::const_iterator::operator*

CUTLASS_HOST_DEVICE const_reference operator*() const

Definition: array_subbyte.h:339

cutlass::Array< T, N, false >::iterator::operator--

CUTLASS_HOST_DEVICE iterator operator--(int)

Definition: array_subbyte.h:247

cutlass::Array< T, N, false >::reference::operator=

CUTLASS_HOST_DEVICE reference & operator=(T x)

Assignment.

Definition: array_subbyte.h:113

cutlass::Array< T, N, false >::end

CUTLASS_HOST_DEVICE iterator end()

Definition: array_subbyte.h:514

cutlass::Array< T, N, false >::reverse_iterator::reverse_iterator

CUTLASS_HOST_DEVICE reverse_iterator(Storage *ptr, int idx=0)

Definition: array_subbyte.h:369

cutlass::Array< T, N, false >::data

CUTLASS_HOST_DEVICE pointer data()

Definition: array_subbyte.h:463

CUTLASS_HOST_DEVICE

#define CUTLASS_HOST_DEVICE

Definition: cutlass.h:89

cutlass::platform::conditional

std::conditional (true specialization)

Definition: platform.h:325

static_assert

#define static_assert(__e, __m)

Definition: platform.h:153

cutlass::Array< T, N, false >::front

CUTLASS_HOST_DEVICE reference front()

Definition: array_subbyte.h:443

[cutlass::Array< T, N, false >::const_reverse_iterator::const_reverse_iterator](classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const reverse iterator.html#a4bef88847b70f6bca81dd46bd883373b)

CUTLASS_HOST_DEVICE const_reverse_iterator(Storage const *ptr, int idx=0)

Definition: array_subbyte.h:389

cutlass::Array< T, N, false >::cbegin

CUTLASS_HOST_DEVICE const_iterator cbegin() const

Definition: array_subbyte.h:509

cutlass::Array< T, N, false >::const_iterator::const_iterator

CUTLASS_HOST_DEVICE const_iterator()

Definition: array_subbyte.h:287

cutlass::Array< T, N, false >::at

CUTLASS_HOST_DEVICE reference at(size_type pos)

Definition: array_subbyte.h:423

CUTLASS_HOST_DEVICE const_reverse_iterator()

Definition: array_subbyte.h:386

cutlass::Array< T, N, false >::const_iterator::operator++

CUTLASS_HOST_DEVICE iterator operator++(int)

Definition: array_subbyte.h:315

cutlass::Array< T, N, false >::raw_data

CUTLASS_HOST_DEVICE Storage const * raw_data() const

Definition: array_subbyte.h:478

cutlass::Array< T, N, false >::const_pointer

value_type const * const_pointer

Definition: array_subbyte.h:87

cutlass::Array< T, N, false >::size

CUTLASS_HOST_DEVICE constexpr size_type size() const

Definition: array_subbyte.h:489

cutlass::Array< T, N, false >::max_size

CUTLASS_HOST_DEVICE constexpr size_type max_size() const

Definition: array_subbyte.h:494

cutlass::Array< T, N, false >::clear

CUTLASS_HOST_DEVICE void clear()

Efficient clear method.

Definition: array_subbyte.h:414

cutlass::Array< T, N, false >::front

CUTLASS_HOST_DEVICE const_reference front() const

Definition: array_subbyte.h:448

cutlass::Array< T, N, false >::iterator::iterator

CUTLASS_HOST_DEVICE iterator()

Definition: array_subbyte.h:208

cutlass::Array< T, N, false >::iterator::iterator

CUTLASS_HOST_DEVICE iterator(Storage *ptr, int idx=0)

Definition: array_subbyte.h:211

cutlass::Array< T, N, false >::Storage

typename platform::conditional< ((kSizeBits%32)!=0), typename platform::conditional< ((kSizeBits%16)!=0), uint8_t, uint16_t >::type, uint32_t >::type Storage

Storage type.

Definition: array_subbyte.h:62

cutlass::Array< T, N, false >::iterator::operator++

CUTLASS_HOST_DEVICE iterator operator++(int)

Definition: array_subbyte.h:236

cutlass::Array< T, N, false >::iterator::operator--

CUTLASS_HOST_DEVICE iterator & operator--()

Definition: array_subbyte.h:224

cutlass::Array< T, N, false >::const_iterator::operator!=

CUTLASS_HOST_DEVICE bool operator!=(iterator const &other) const

Definition: array_subbyte.h:349

cutlass::Array< T, N, false >::iterator::operator!=