docs/subbyte__reference_8h_source.html
| | 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
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
97ConstSubbyteReference(): ptr_(nullptr), offset_(0) { }
98
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
118Element *ptr = nullptr
119 ): ConstSubbyteReference(ptr, 0) { }
120
123StoragePointer storage_pointer() const {
124return ptr_;
125 }
126
129int element_offset() const {
130return offset_;
131 }
132
136Storage item = Storage((*ptr_ >> (offset_ * sizeof_bits<Element>::value)) & kMask);
137return reinterpret_cast<Element const &>(item);
138 }
139
143return get();
144 }
145
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
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
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
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
214ConstSubbyteReference operator+(int offset) const {
215
216ConstSubbyteReference ref(ptr_, offset_);
217 ref += offset;
218
219return ref;
220 }
221
224ConstSubbyteReference operator+(long long offset) const {
225
226ConstSubbyteReference ref(ptr_, offset_);
227 ref += offset;
228
229return ref;
230 }
231
234ConstSubbyteReference operator-(int offset) const {
235
236ConstSubbyteReference ref(ptr_, offset_);
237 ref -= offset;
238
239return ref;
240 }
241
244ConstSubbyteReference operator-=(long long offset) const {
245
246ConstSubbyteReference ref(ptr_, offset_);
247 ref -= offset;
248
249return ref;
250 }
251
254 ptrdiff_t operator-(ConstSubbyteReference ref) const {
255return (ptr_ - ref.ptr_) * kElementsPerVector + (offset_ - ref.offset_);
256 }
257
260explicit operator int() const {
261return int(get());
262 }
263
266explicit operator int64_t() const {
267return int64_t(get());
268 }
269
272explicit operator uint64_t() const {
273return uint64_t(get());
274 }
275
278explicit operator float() const {
279return float(get());
280 }
281
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
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
331SubbyteReference(): ptr_(nullptr), offset_(0) { }
332
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
352Element *ptr = nullptr
353 ): SubbyteReference(ptr, 0) { }
354
357StoragePointer storage_pointer() const {
358return ptr_;
359 }
360
363int element_offset() const {
364return offset_;
365 }
366
370Storage item = Storage((*ptr_ >> (offset_ * sizeof_bits<Element>::value)) & kMask);
371return reinterpret_cast<Element const &>(item);
372 }
373
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
389return get();
390 }
391
394SubbyteReference &operator=(Element const & x) {
395return set(x);
396 }
397
400SubbyteReference &operator=(SubbyteReference const & x) {
401return set(x.get());
402 }
403
406SubbyteReference &operator=(
407ConstSubbyteReference<Element, Storage> const &x) {
408return set(x.get());
409 }
410
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
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
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
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
479SubbyteReference operator+(int offset) const {
480
481SubbyteReference ref(ptr_, offset_);
482 ref += offset;
483
484return ref;
485 }
486
489SubbyteReference operator+(long long offset) const {
490
491SubbyteReference ref(ptr_, offset_);
492 ref += offset;
493
494return ref;
495 }
496
499SubbyteReference operator-(int offset) const {
500
501SubbyteReference ref(ptr_, offset_);
502 ref -= offset;
503
504return ref;
505 }
506
509SubbyteReference operator-=(long long offset) const {
510
511SubbyteReference ref(ptr_, offset_);
512 ref -= offset;
513
514return ref;
515 }
516
519 ptrdiff_t operator-(SubbyteReference ref) const {
520return (ptr_ - ref.ptr_) * kElementsPerVector + (offset_ - ref.offset_);
521 }
522
525explicit operator int() const {
526return int(get());
527 }
528
531explicit operator int64_t() const {
532return int64_t(get());
533 }
534
537explicit operator uint64_t() const {
538return uint64_t(get());
539 }
540
543explicit operator float() const {
544return float(get());
545 }
546
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> {
562static Element &get(Element *ptr, int64_t offset) {
563return ptr[offset];
564 }
565
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> {
575static SubbyteReference<Element> get(Element *ptr, int64_t offset) {
576return SubbyteReference<Element>(ptr, offset);
577 }
578
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
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
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
#define nullptr
nullptr
Definition: platform.h:144
cutlass::SubbyteReference::Element
Element_ Element
Definition: subbyte_reference.h:297
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
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:89
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
#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
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