CUTLASS
CUDA Templates for Linear Algebra Subroutines and Solvers
functional.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  **************************************************************************************************/
31 #pragma once
32 
33 #include "cutlass/cutlass.h"
34 #include "cutlass/numeric_types.h"
35 
36 #include "cutlass/complex.h"
37 
38 #include "cutlass/array.h"
39 #include "cutlass/half.h"
40 
41 namespace cutlass {
42 
44 
45 template <typename T>
46 struct plus {
47  CUTLASS_HOST_DEVICE
48  T operator()(T lhs, T const &rhs) const {
49  lhs += rhs;
50  return lhs;
51  }
52 };
53 
54 template <typename T>
55 struct minus {
56  CUTLASS_HOST_DEVICE
57  T operator()(T lhs, T const &rhs) const {
58  lhs -= rhs;
59  return lhs;
60  }
61 };
62 
63 template <typename T>
64 struct multiplies {
65  CUTLASS_HOST_DEVICE
66  T operator()(T lhs, T const &rhs) const {
67  lhs *= rhs;
68  return lhs;
69  }
70 };
71 
72 template <typename T>
73 struct divides {
74  CUTLASS_HOST_DEVICE
75  T operator()(T lhs, T const &rhs) const {
76  lhs /= rhs;
77  return lhs;
78  }
79 };
80 
81 
82 template <typename T>
83 struct negate {
84  CUTLASS_HOST_DEVICE
85  T operator()(T lhs) const {
86  return -lhs;
87  }
88 };
89 
91 template <typename A, typename B = A, typename C = A>
92 struct multiply_add {
93  CUTLASS_HOST_DEVICE
94  C operator()(A const &a, B const &b, C const &c) const {
95  return C(a) * C(b) + c;
96  }
97 };
98 
100 template <typename T>
101 struct xor_add {
102  CUTLASS_HOST_DEVICE
103  T operator()(T const &a, T const &b, T const &c) const {
104  return ((a ^ b) + c);
105  }
106 };
107 
109 //
110 // Partial specialization for complex<T> to target four scalar fused multiply-adds.
111 //
113 
115 template <typename T>
116 struct multiply_add<complex<T>, complex<T>, complex<T>> {
117  CUTLASS_HOST_DEVICE
118  complex<T> operator()(
119  complex<T> const &a,
120  complex<T> const &b,
121  complex<T> const &c) const {
122 
123  T real = c.real();
124  T imag = c.imag();
125 
126  real += a.real() * b.real();
127  real += -a.imag() * b.imag();
128  imag += a.real() * b.imag();
129  imag += a.imag () * b.real();
130 
131  return complex<T>{
132  real,
133  imag
134  };
135  }
136 };
137 
139 template <typename T>
140 struct multiply_add<complex<T>, T, complex<T>> {
141  CUTLASS_HOST_DEVICE
142  complex<T> operator()(
143  complex<T> const &a,
144  T const &b,
145  complex<T> const &c) const {
146 
147  T real = c.real();
148  T imag = c.imag();
149 
150  real += a.real() * b;
151  imag += a.imag () * b;
152 
153  return complex<T>{
154  real,
155  imag
156  };
157  }
158 };
159 
161 template <typename T>
162 struct multiply_add<T, complex<T>, complex<T>> {
163  CUTLASS_HOST_DEVICE
164  complex<T> operator()(
165  T const &a,
166  complex<T> const &b,
167  complex<T> const &c) const {
168 
169  T real = c.real();
170  T imag = c.imag();
171 
172  real += a * b.real();
173  imag += a * b.imag();
174 
175  return complex<T>{
176  real,
177  imag
178  };
179  }
180 };
181 
183 //
184 // Partial specializations for Array<T, N>
185 //
187 
188 template <typename T, int N>
189 struct plus<Array<T, N>> {
190  CUTLASS_HOST_DEVICE
191  Array<T, N> operator()(Array<T, N> const &lhs, Array<T, N> const &rhs) const {
192 
193  Array<T, N> result;
194  plus<T> scalar_op;
195 
196  CUTLASS_PRAGMA_UNROLL
197  for (int i = 0; i < N; ++i) {
198  result[i] = scalar_op(lhs[i], rhs[i]);
199  }
200 
201  return result;
202  }
203 
204  CUTLASS_HOST_DEVICE
205  Array<T, N> operator()(Array<T, N> const &lhs, T const &scalar) const {
206 
207  Array<T, N> result;
208  plus<T> scalar_op;
209 
210  CUTLASS_PRAGMA_UNROLL
211  for (int i = 0; i < N; ++i) {
212  result[i] = scalar_op(lhs[i], scalar);
213  }
214 
215  return result;
216  }
217 
218  CUTLASS_HOST_DEVICE
219  Array<T, N> operator()( T const &scalar, Array<T, N> const &rhs) const {
220 
221  Array<T, N> result;
222  plus<T> scalar_op;
223 
224  CUTLASS_PRAGMA_UNROLL
225  for (int i = 0; i < N; ++i) {
226  result[i] = scalar_op(scalar, rhs[i]);
227  }
228 
229  return result;
230  }
231 };
232 
233 
234 template <typename T>
235 struct maximum {
236 
237  CUTLASS_HOST_DEVICE
238  T operator()(T const &lhs, T const &rhs) const {
239  return (lhs < rhs ? rhs : lhs);
240  }
241 };
242 
243 template <>
244 struct maximum<float> {
245  CUTLASS_HOST_DEVICE
246  float operator()(float const &lhs, float const &rhs) const {
247  return fmaxf(lhs, rhs);
248  }
249 };
250 
251 template <typename T, int N>
252 struct maximum<Array<T, N>> {
253 
254  CUTLASS_HOST_DEVICE
255  Array<T, N> operator()(Array<T, N> const &lhs, Array<T, N> const &rhs) const {
256 
257  Array<T, N> result;
258  maximum<T> scalar_op;
259 
260  CUTLASS_PRAGMA_UNROLL
261  for (int i = 0; i < N; ++i) {
262  result[i] = scalar_op(lhs[i], rhs[i]);
263  }
264 
265  return result;
266  }
267 
268  CUTLASS_HOST_DEVICE
269  Array<T, N> operator()(Array<T, N> const &lhs, T const &scalar) const {
270 
271  Array<T, N> result;
272  maximum<T> scalar_op;
273 
274  CUTLASS_PRAGMA_UNROLL
275  for (int i = 0; i < N; ++i) {
276  result[i] = scalar_op(lhs[i], scalar);
277  }
278 
279  return result;
280  }
281 
282  CUTLASS_HOST_DEVICE
283  Array<T, N> operator()( T const &scalar, Array<T, N> const &rhs) const {
284 
285  Array<T, N> result;
286  maximum<T> scalar_op;
287 
288  CUTLASS_PRAGMA_UNROLL
289  for (int i = 0; i < N; ++i) {
290  result[i] = scalar_op(scalar, rhs[i]);
291  }
292 
293  return result;
294  }
295 };
296 
297 template <typename T>
298 struct minimum {
299 
300  CUTLASS_HOST_DEVICE
301  T operator()(T const &lhs, T const &rhs) const {
302  return (rhs < lhs ? rhs : lhs);
303  }
304 };
305 
306 template <>
307 struct minimum<float> {
308  CUTLASS_HOST_DEVICE
309  float operator()(float const &lhs, float const &rhs) const {
310  return fminf(lhs, rhs);
311  }
312 };
313 
314 template <typename T, int N>
315 struct minimum<Array<T, N>> {
316 
317  CUTLASS_HOST_DEVICE
318  static T scalar_op(T const &lhs, T const &rhs) {
319  return (rhs < lhs ? rhs : lhs);
320  }
321 
322  CUTLASS_HOST_DEVICE
323  Array<T, N> operator()(Array<T, N> const &lhs, Array<T, N> const &rhs) const {
324 
325  Array<T, N> result;
326  minimum<T> scalar_op;
327 
328  CUTLASS_PRAGMA_UNROLL
329  for (int i = 0; i < N; ++i) {
330  result[i] = scalar_op(lhs[i], rhs[i]);
331  }
332 
333  return result;
334  }
335 
336  CUTLASS_HOST_DEVICE
337  Array<T, N> operator()(Array<T, N> const &lhs, T const &scalar) const {
338 
339  Array<T, N> result;
340  minimum<T> scalar_op;
341 
342  CUTLASS_PRAGMA_UNROLL
343  for (int i = 0; i < N; ++i) {
344  result[i] = scalar_op(lhs[i], scalar);
345  }
346 
347  return result;
348  }
349 
350  CUTLASS_HOST_DEVICE
351  Array<T, N> operator()( T const &scalar, Array<T, N> const &rhs) const {
352 
353  Array<T, N> result;
354  minimum<T> scalar_op;
355 
356  CUTLASS_PRAGMA_UNROLL
357  for (int i = 0; i < N; ++i) {
358  result[i] = scalar_op(scalar, rhs[i]);
359  }
360 
361  return result;
362  }
363 };
364 
365 template <typename T, int N>
366 struct minus<Array<T, N>> {
367 
368  CUTLASS_HOST_DEVICE
369  Array<T, N> operator()(Array<T, N> const &lhs, Array<T, N> const &rhs) const {
370 
371  Array<T, N> result;
372  minus<T> scalar_op;
373 
374  CUTLASS_PRAGMA_UNROLL
375  for (int i = 0; i < N; ++i) {
376  result[i] = scalar_op(lhs[i], rhs[i]);
377  }
378 
379  return result;
380  }
381 
382  CUTLASS_HOST_DEVICE
383  Array<T, N> operator()(Array<T, N> const &lhs, T const &scalar) const {
384 
385  Array<T, N> result;
386  minus<T> scalar_op;
387 
388  CUTLASS_PRAGMA_UNROLL
389  for (int i = 0; i < N; ++i) {
390  result[i] = scalar_op(lhs[i], scalar);
391  }
392 
393  return result;
394  }
395 
396  CUTLASS_HOST_DEVICE
397  Array<T, N> operator()( T const &scalar, Array<T, N> const &rhs) const {
398 
399  Array<T, N> result;
400  minus<T> scalar_op;
401 
402  CUTLASS_PRAGMA_UNROLL
403  for (int i = 0; i < N; ++i) {
404  result[i] = scalar_op(scalar, rhs[i]);
405  }
406 
407  return result;
408  }
409 };
410 
411 template <typename T, int N>
412 struct multiplies<Array<T, N>> {
413 
414  CUTLASS_HOST_DEVICE
415  Array<T, N> operator()(Array<T, N> const &lhs, Array<T, N> const &rhs) const {
416 
417  Array<T, N> result;
418  multiplies<T> scalar_op;
419 
420  CUTLASS_PRAGMA_UNROLL
421  for (int i = 0; i < N; ++i) {
422  result[i] = scalar_op(lhs[i], rhs[i]);
423  }
424 
425  return result;
426  }
427 
428  CUTLASS_HOST_DEVICE
429  Array<T, N> operator()(Array<T, N> const &lhs, T const &scalar) const {
430 
431  Array<T, N> result;
432  multiplies<T> scalar_op;
433 
434  CUTLASS_PRAGMA_UNROLL
435  for (int i = 0; i < N; ++i) {
436  result[i] = scalar_op(lhs[i], scalar);
437  }
438 
439  return result;
440  }
441 
442  CUTLASS_HOST_DEVICE
443  Array<T, N> operator()( T const &scalar, Array<T, N> const &rhs) const {
444 
445  Array<T, N> result;
446  multiplies<T> scalar_op;
447 
448  CUTLASS_PRAGMA_UNROLL
449  for (int i = 0; i < N; ++i) {
450  result[i] = scalar_op(scalar, rhs[i]);
451  }
452 
453  return result;
454  }
455 };
456 
457 template <typename T, int N>
458 struct divides<Array<T, N>> {
459 
460  CUTLASS_HOST_DEVICE
461  Array<T, N> operator()(Array<T, N> const &lhs, Array<T, N> const &rhs) const {
462 
463  Array<T, N> result;
464  divides<T> scalar_op;
465 
466  CUTLASS_PRAGMA_UNROLL
467  for (int i = 0; i < N; ++i) {
468  result[i] = scalar_op(lhs[i], rhs[i]);
469  }
470 
471  return result;
472  }
473 
474  CUTLASS_HOST_DEVICE
475  Array<T, N> operator()(Array<T, N> const &lhs, T const &scalar) const {
476 
477  Array<T, N> result;
478  divides<T> scalar_op;
479 
480  CUTLASS_PRAGMA_UNROLL
481  for (int i = 0; i < N; ++i) {
482  result[i] = scalar_op(lhs[i], scalar);
483  }
484 
485  return result;
486  }
487 
488  CUTLASS_HOST_DEVICE
489  Array<T, N> operator()( T const &scalar, Array<T, N> const &rhs) const {
490 
491  Array<T, N> result;
492  divides<T> scalar_op;
493 
494  CUTLASS_PRAGMA_UNROLL
495  for (int i = 0; i < N; ++i) {
496  result[i] = scalar_op(scalar, rhs[i]);
497  }
498 
499  return result;
500  }
501 };
502 
503 
504 template <typename T, int N>
505 struct negate<Array<T, N>> {
506 
507  CUTLASS_HOST_DEVICE
508  Array<T, N> operator()(Array<T, N> const &lhs) const {
509 
510  Array<T, N> result;
511  negate<T> scalar_op;
512 
513  CUTLASS_PRAGMA_UNROLL
514  for (int i = 0; i < N; ++i) {
515  result[i] = scalar_op(lhs[i]);
516  }
517 
518  return result;
519  }
520 };
521 
523 template <typename T, int N>
524 struct multiply_add<Array<T, N>, Array<T, N>, Array<T, N>> {
525 
526  CUTLASS_HOST_DEVICE
527  Array<T, N> operator()(Array<T, N> const &a, Array<T, N> const &b, Array<T, N> const &c) const {
528 
529  Array<T, N> result;
530  multiply_add<T> scalar_op;
531 
532  CUTLASS_PRAGMA_UNROLL
533  for (int i = 0; i < N; ++i) {
534  result[i] = scalar_op(a[i], b[i], c[i]);
535  }
536 
537  return result;
538  }
539 
540  CUTLASS_HOST_DEVICE
541  Array<T, N> operator()(Array<T, N> const &a, T const &scalar, Array<T, N> const &c) const {
542 
543  Array<T, N> result;
544  multiply_add<T> scalar_op;
545 
546  CUTLASS_PRAGMA_UNROLL
547  for (int i = 0; i < N; ++i) {
548  result[i] = scalar_op(a[i], scalar, c[i]);
549  }
550 
551  return result;
552  }
553 
554  CUTLASS_HOST_DEVICE
555  Array<T, N> operator()(T const &scalar, Array<T, N> const &b, Array<T, N> const &c) const {
556 
557  Array<T, N> result;
558  multiply_add<T> scalar_op;
559 
560  CUTLASS_PRAGMA_UNROLL
561  for (int i = 0; i < N; ++i) {
562  result[i] = scalar_op(scalar, b[i], c[i]);
563  }
564 
565  return result;
566  }
567 };
568 
570 //
571 // Partial specializations for Array<half_t, N> targeting SIMD instructions in device code.
572 //
574 
575 template <int N>
576 struct plus<Array<half_t, N>> {
577  CUTLASS_HOST_DEVICE
578  Array<half_t, N> operator()(Array<half_t, N> const & lhs, Array<half_t, N> const &rhs) const {
579  Array<half_t, N> result;
580  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)
581 
582  __half2 *result_ptr = reinterpret_cast<__half2 *>(&result);
583  __half2 const *lhs_ptr = reinterpret_cast<__half2 const *>(&lhs);
584  __half2 const *rhs_ptr = reinterpret_cast<__half2 const *>(&rhs);
585 
586  CUTLASS_PRAGMA_UNROLL
587  for (int i = 0; i < N / 2; ++i) {
588  result_ptr[i] = __hadd2(lhs_ptr[i], rhs_ptr[i]);
589  }
590 
591  if (N % 2) {
592  __half const *a_residual_ptr = reinterpret_cast<__half const *>(&lhs);
593  __half const *b_residual_ptr = reinterpret_cast<__half const *>(&rhs);
594  __half d_residual = __hadd(a_residual_ptr[N - 1], b_residual_ptr[N - 1]);
595 
596  result[N - 1] = reinterpret_cast<half_t const &>(d_residual);
597  }
598 
599  #else
600 
601  CUTLASS_PRAGMA_UNROLL
602  for (int i = 0; i < N; ++i) {
603  result[i] = lhs[i] + rhs[i];
604  }
605  #endif
606 
607  return result;
608  }
609 
610  CUTLASS_HOST_DEVICE
611  Array<half_t, N> operator()(half_t const & lhs, Array<half_t, N> const &rhs) const {
612  Array<half_t, N> result;
613  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)
614 
615  __half2 *result_ptr = reinterpret_cast<__half2 *>(&result);
616  __half2 lhs_pair = __half2half2(reinterpret_cast<__half const &>(lhs));
617  __half2 const *rhs_ptr = reinterpret_cast<__half2 const *>(&rhs);
618 
619  CUTLASS_PRAGMA_UNROLL
620  for (int i = 0; i < N / 2; ++i) {
621  result_ptr[i] = __hadd2(lhs_pair, rhs_ptr[i]);
622  }
623 
624  if (N % 2) {
625  __half const *b_residual_ptr = reinterpret_cast<__half const *>(&rhs);
626  __half d_residual = __hadd(reinterpret_cast<__half const &>(lhs), b_residual_ptr[N - 1]);
627 
628  result[N - 1] = reinterpret_cast<half_t const &>(d_residual);
629  }
630 
631  #else
632 
633  CUTLASS_PRAGMA_UNROLL
634  for (int i = 0; i < N; ++i) {
635  result[i] = lhs + rhs[i];
636  }
637  #endif
638 
639  return result;
640  }
641 
642  CUTLASS_HOST_DEVICE
643  Array<half_t, N> operator()(Array<half_t, N> const & lhs, half_t const &rhs) const {
644  Array<half_t, N> result;
645  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)
646 
647  __half2 *result_ptr = reinterpret_cast<__half2 *>(&result);
648  __half2 const *lhs_ptr = reinterpret_cast<__half2 const *>(&lhs);
649  __half2 rhs_pair = __half2half2(reinterpret_cast<__half const &>(rhs));
650 
651  CUTLASS_PRAGMA_UNROLL
652  for (int i = 0; i < N / 2; ++i) {
653  result_ptr[i] = __hadd2(lhs_ptr[i], rhs_pair);
654  }
655 
656  if (N % 2) {
657  __half const *a_residual_ptr = reinterpret_cast<__half const *>(&lhs);
658  __half d_residual = __hadd(a_residual_ptr[N - 1], reinterpret_cast<__half const &>(rhs));
659 
660  result[N - 1] = reinterpret_cast<half_t const &>(d_residual);
661  }
662 
663  #else
664 
665  CUTLASS_PRAGMA_UNROLL
666  for (int i = 0; i < N; ++i) {
667  result[i] = lhs[i] + rhs;
668  }
669  #endif
670 
671  return result;
672  }
673 };
674 
675 template <int N>
676 struct minus<Array<half_t, N>> {
677  CUTLASS_HOST_DEVICE
678  Array<half_t, N> operator()(Array<half_t, N> const & lhs, Array<half_t, N> const &rhs) const {
679  Array<half_t, N> result;
680  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)
681 
682  __half2 *result_ptr = reinterpret_cast<__half2 *>(&result);
683  __half2 const *lhs_ptr = reinterpret_cast<__half2 const *>(&lhs);
684  __half2 const *rhs_ptr = reinterpret_cast<__half2 const *>(&rhs);
685 
686  CUTLASS_PRAGMA_UNROLL
687  for (int i = 0; i < N / 2; ++i) {
688  result_ptr[i] = __hsub2(lhs_ptr[i], rhs_ptr[i]);
689  }
690 
691  if (N % 2) {
692  __half const *a_residual_ptr = reinterpret_cast<__half const *>(&lhs);
693  __half const *b_residual_ptr = reinterpret_cast<__half const *>(&rhs);
694  __half d_residual = __hsub(a_residual_ptr[N - 1], b_residual_ptr[N - 1]);
695 
696  result[N - 1] = reinterpret_cast<half_t const &>(d_residual);
697  }
698 
699  #else
700 
701  CUTLASS_PRAGMA_UNROLL
702  for (int i = 0; i < N; ++i) {
703  result[i] = lhs[i] - rhs[i];
704  }
705  #endif
706 
707  return result;
708  }
709 
710  CUTLASS_HOST_DEVICE
711  Array<half_t, N> operator()(half_t const & lhs, Array<half_t, N> const &rhs) const {
712  Array<half_t, N> result;
713  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)
714 
715  __half2 *result_ptr = reinterpret_cast<__half2 *>(&result);
716  __half2 lhs_pair = __half2half2(reinterpret_cast<__half const &>(lhs));
717  __half2 const *rhs_ptr = reinterpret_cast<__half2 const *>(&rhs);
718 
719  CUTLASS_PRAGMA_UNROLL
720  for (int i = 0; i < N / 2; ++i) {
721  result_ptr[i] = __hsub2(lhs_pair, rhs_ptr[i]);
722  }
723 
724  if (N % 2) {
725  __half const *b_residual_ptr = reinterpret_cast<__half const *>(&rhs);
726  __half d_residual = __hsub(reinterpret_cast<__half const &>(lhs), b_residual_ptr[N - 1]);
727 
728  result[N - 1] = reinterpret_cast<half_t const &>(d_residual);
729  }
730 
731  #else
732 
733  CUTLASS_PRAGMA_UNROLL
734  for (int i = 0; i < N; ++i) {
735  result[i] = lhs - rhs[i];
736  }
737  #endif
738 
739  return result;
740  }
741 
742  CUTLASS_HOST_DEVICE
743  Array<half_t, N> operator()(Array<half_t, N> const & lhs, half_t const &rhs) const {
744  Array<half_t, N> result;
745  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)
746 
747  __half2 *result_ptr = reinterpret_cast<__half2 *>(&result);
748  __half2 const *lhs_ptr = reinterpret_cast<__half2 const *>(&lhs);
749  __half2 rhs_pair = __half2half2(reinterpret_cast<__half const &>(rhs));
750 
751  CUTLASS_PRAGMA_UNROLL
752  for (int i = 0; i < N / 2; ++i) {
753  result_ptr[i] = __hsub2(lhs_ptr[i], rhs_pair);
754  }
755 
756  if (N % 2) {
757  __half const *a_residual_ptr = reinterpret_cast<__half const *>(&lhs);
758  __half d_residual = __hsub(a_residual_ptr[N - 1], reinterpret_cast<__half const &>(rhs));
759 
760  result[N - 1] = reinterpret_cast<half_t const &>(d_residual);
761  }
762 
763  #else
764 
765  CUTLASS_PRAGMA_UNROLL
766  for (int i = 0; i < N; ++i) {
767  result[i] = lhs[i] - rhs;
768  }
769  #endif
770 
771  return result;
772  }
773 };
774 
775 template <int N>
776 struct multiplies<Array<half_t, N>> {
777  CUTLASS_HOST_DEVICE
778  Array<half_t, N> operator()(Array<half_t, N> const & lhs, Array<half_t, N> const &rhs) const {
779  Array<half_t, N> result;
780  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)
781 
782  __half2 *result_ptr = reinterpret_cast<__half2 *>(&result);
783  __half2 const *lhs_ptr = reinterpret_cast<__half2 const *>(&lhs);
784  __half2 const *rhs_ptr = reinterpret_cast<__half2 const *>(&rhs);
785 
786  CUTLASS_PRAGMA_UNROLL
787  for (int i = 0; i < N / 2; ++i) {
788  result_ptr[i] = __hmul2(lhs_ptr[i], rhs_ptr[i]);
789  }
790 
791  if (N % 2) {
792  __half const *a_residual_ptr = reinterpret_cast<__half const *>(&lhs);
793  __half const *b_residual_ptr = reinterpret_cast<__half const *>(&rhs);
794  __half d_residual = __hmul(a_residual_ptr[N - 1], b_residual_ptr[N - 1]);
795 
796  result[N - 1] = reinterpret_cast<half_t const &>(d_residual);
797  }
798 
799  #else
800 
801  CUTLASS_PRAGMA_UNROLL
802  for (int i = 0; i < N; ++i) {
803  result[i] = lhs[i] * rhs[i];
804  }
805  #endif
806 
807  return result;
808  }
809 
810  CUTLASS_HOST_DEVICE
811  Array<half_t, N> operator()(half_t const & lhs, Array<half_t, N> const &rhs) const {
812  Array<half_t, N> result;
813  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)
814 
815  __half2 *result_ptr = reinterpret_cast<__half2 *>(&result);
816  __half2 lhs_pair = __half2half2(reinterpret_cast<__half const &>(lhs));
817  __half2 const *rhs_ptr = reinterpret_cast<__half2 const *>(&rhs);
818 
819  CUTLASS_PRAGMA_UNROLL
820  for (int i = 0; i < N / 2; ++i) {
821  result_ptr[i] = __hmul2(lhs_pair, rhs_ptr[i]);
822  }
823 
824  if (N % 2) {
825  __half const *b_residual_ptr = reinterpret_cast<__half const *>(&rhs);
826 
827  __half d_residual = __hmul(
828  reinterpret_cast<__half const &>(lhs),
829  b_residual_ptr[N - 1]);
830 
831  result[N - 1] = reinterpret_cast<half_t const &>(d_residual);
832  }
833 
834  #else
835 
836  CUTLASS_PRAGMA_UNROLL
837  for (int i = 0; i < N; ++i) {
838  result[i] = lhs * rhs[i];
839  }
840  #endif
841 
842  return result;
843  }
844 
845  CUTLASS_HOST_DEVICE
846  Array<half_t, N> operator()(Array<half_t, N> const & lhs, half_t const &rhs) const {
847  Array<half_t, N> result;
848  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)
849 
850  __half2 *result_ptr = reinterpret_cast<__half2 *>(&result);
851  __half2 const *lhs_ptr = reinterpret_cast<__half2 const *>(&lhs);
852  __half2 rhs_pair = __half2half2(reinterpret_cast<__half const &>(rhs));
853 
854  CUTLASS_PRAGMA_UNROLL
855  for (int i = 0; i < N / 2; ++i) {
856  result_ptr[i] = __hmul2(lhs_ptr[i], rhs_pair);
857  }
858 
859  if (N % 2) {
860  __half const *a_residual_ptr = reinterpret_cast<__half const *>(&lhs);
861 
862  __half d_residual = __hmul(
863  a_residual_ptr[N - 1],
864  reinterpret_cast<__half const &>(rhs));
865 
866  result[N - 1] = reinterpret_cast<half_t const &>(d_residual);
867  }
868 
869  #else
870 
871  CUTLASS_PRAGMA_UNROLL
872  for (int i = 0; i < N; ++i) {
873  result[i] = lhs[i] * rhs;
874  }
875  #endif
876 
877  return result;
878  }
879 };
880 
881 template <int N>
882 struct divides<Array<half_t, N>> {
883  CUTLASS_HOST_DEVICE
884  Array<half_t, N> operator()(Array<half_t, N> const & lhs, Array<half_t, N> const &rhs) const {
885  Array<half_t, N> result;
886  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)
887 
888  __half2 *result_ptr = reinterpret_cast<__half2 *>(&result);
889  __half2 const *lhs_ptr = reinterpret_cast<__half2 const *>(&lhs);
890  __half2 const *rhs_ptr = reinterpret_cast<__half2 const *>(&rhs);
891 
892  CUTLASS_PRAGMA_UNROLL
893  for (int i = 0; i < N / 2; ++i) {
894  result_ptr[i] = __h2div(lhs_ptr[i], rhs_ptr[i]);
895  }
896 
897  if (N % 2) {
898  __half const *a_residual_ptr = reinterpret_cast<__half const *>(&lhs);
899  __half const *b_residual_ptr = reinterpret_cast<__half const *>(&rhs);
900 
901  __half d_residual = __hdiv(
902  a_residual_ptr[N - 1],
903  b_residual_ptr[N - 1]);
904 
905  result[N - 1] = reinterpret_cast<half_t const &>(d_residual);
906  }
907 
908  #else
909 
910  CUTLASS_PRAGMA_UNROLL
911  for (int i = 0; i < N; ++i) {
912  result[i] = lhs[i] / rhs[i];
913  }
914  #endif
915 
916  return result;
917  }
918 
919  CUTLASS_HOST_DEVICE
920  Array<half_t, N> operator()(half_t const & lhs, Array<half_t, N> const &rhs) const {
921  Array<half_t, N> result;
922  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)
923 
924  __half2 *result_ptr = reinterpret_cast<__half2 *>(&result);
925  __half2 lhs_pair = __half2half2(reinterpret_cast<__half const &>(lhs));
926  __half2 const *rhs_ptr = reinterpret_cast<__half2 const *>(&rhs);
927 
928  CUTLASS_PRAGMA_UNROLL
929  for (int i = 0; i < N / 2; ++i) {
930  result_ptr[i] = __h2div(lhs_pair, rhs_ptr[i]);
931  }
932 
933  if (N % 2) {
934  __half const *b_residual_ptr = reinterpret_cast<__half const *>(&rhs);
935 
936  __half d_residual = __hdiv(
937  reinterpret_cast<__half const &>(lhs),
938  b_residual_ptr[N - 1]);
939 
940  result[N - 1] = reinterpret_cast<half_t const &>(d_residual);
941  }
942 
943  #else
944 
945  CUTLASS_PRAGMA_UNROLL
946  for (int i = 0; i < N; ++i) {
947  result[i] = lhs / rhs[i];
948  }
949  #endif
950 
951  return result;
952  }
953 
954  CUTLASS_HOST_DEVICE
955  Array<half_t, N> operator()(Array<half_t, N> const & lhs, half_t const &rhs) const {
956  Array<half_t, N> result;
957  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)
958 
959  __half2 *result_ptr = reinterpret_cast<__half2 *>(&result);
960  __half2 const *lhs_ptr = reinterpret_cast<__half2 const *>(&lhs);
961  __half2 rhs_pair = __half2half2(reinterpret_cast<__half const &>(rhs));
962 
963  CUTLASS_PRAGMA_UNROLL
964  for (int i = 0; i < N / 2; ++i) {
965  result_ptr[i] = __h2div(lhs_ptr[i], rhs_pair);
966  }
967 
968  if (N % 2) {
969  __half const *a_residual_ptr = reinterpret_cast<__half const *>(&lhs);
970 
971  __half d_residual = __hdiv(
972  a_residual_ptr[N - 1],
973  reinterpret_cast<__half const &>(rhs));
974 
975  result[N - 1] = reinterpret_cast<half_t const &>(d_residual);
976  }
977 
978  #else
979 
980  CUTLASS_PRAGMA_UNROLL
981  for (int i = 0; i < N; ++i) {
982  result[i] = lhs[i] / rhs;
983  }
984  #endif
985 
986  return result;
987  }
988 };
989 
990 template <int N>
991 struct negate<Array<half_t, N>> {
992  CUTLASS_HOST_DEVICE
993  Array<half_t, N> operator()(Array<half_t, N> const & lhs) const {
994  Array<half_t, N> result;
995  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)
996 
997  __half2 *result_ptr = reinterpret_cast<__half2 *>(&result);
998  __half2 const *source_ptr = reinterpret_cast<__half2 const *>(&lhs);
999 
1000  CUTLASS_PRAGMA_UNROLL
1001  for (int i = 0; i < N / 2; ++i) {
1002  result_ptr[i] = __hneg2(source_ptr[i]);
1003  }
1004 
1005  if (N % 2) {
1006  half_t x = lhs[N - 1];
1007  __half lhs_val = -reinterpret_cast<__half const &>(x);
1008  result[N - 1] = reinterpret_cast<half_t const &>(lhs_val);
1009  }
1010 
1011  #else
1012 
1013  CUTLASS_PRAGMA_UNROLL
1014  for (int i = 0; i < N; ++i) {
1015  result[i] = -lhs[i];
1016  }
1017  #endif
1018 
1019  return result;
1020  }
1021 };
1022 
1024 template <int N>
1025 struct multiply_add<Array<half_t, N>, Array<half_t, N>, Array<half_t, N>> {
1026 
1027  CUTLASS_HOST_DEVICE
1028  Array<half_t, N> operator()(
1029  Array<half_t, N> const &a,
1030  Array<half_t, N> const &b,
1031  Array<half_t, N> const &c) const {
1032 
1033  Array<half_t, N> result;
1034  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)
1035 
1036  __half2 *result_ptr = reinterpret_cast<__half2 *>(&result);
1037  __half2 const *a_ptr = reinterpret_cast<__half2 const *>(&a);
1038  __half2 const *b_ptr = reinterpret_cast<__half2 const *>(&b);
1039  __half2 const *c_ptr = reinterpret_cast<__half2 const *>(&c);
1040 
1041  CUTLASS_PRAGMA_UNROLL
1042  for (int i = 0; i < N / 2; ++i) {
1043  result_ptr[i] = __hfma2(a_ptr[i], b_ptr[i], c_ptr[i]);
1044  }
1045 
1046  if (N % 2) {
1047 
1048  __half const *a_residual_ptr = reinterpret_cast<__half const *>(&a);
1049  __half const *b_residual_ptr = reinterpret_cast<__half const *>(&b);
1050  __half const *c_residual_ptr = reinterpret_cast<__half const *>(&c);
1051 
1052  __half d_residual = __hfma(
1053  a_residual_ptr[N - 1],
1054  b_residual_ptr[N - 1],
1055  c_residual_ptr[N - 1]);
1056 
1057  result[N - 1] = reinterpret_cast<half_t const &>(d_residual);
1058  }
1059 
1060  #else
1061 
1062  multiply_add<half_t> op;
1063 
1064  CUTLASS_PRAGMA_UNROLL
1065  for (int i = 0; i < N; ++i) {
1066  result[i] = op(a[i], b[i], c[i]);
1067  }
1068  #endif
1069 
1070  return result;
1071  }
1072 
1073  CUTLASS_HOST_DEVICE
1074  Array<half_t, N> operator()(
1075  half_t const &a,
1076  Array<half_t, N> const &b,
1077  Array<half_t, N> const &c) const {
1078 
1079  Array<half_t, N> result;
1080  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)
1081 
1082  __half2 *result_ptr = reinterpret_cast<__half2 *>(&result);
1083  __half2 a_pair = __half2half2(reinterpret_cast<__half const &>(a));
1084  __half2 const *b_ptr = reinterpret_cast<__half2 const *>(&b);
1085  __half2 const *c_ptr = reinterpret_cast<__half2 const *>(&c);
1086 
1087  CUTLASS_PRAGMA_UNROLL
1088  for (int i = 0; i < N / 2; ++i) {
1089  result_ptr[i] = __hfma2(a_pair, b_ptr[i], c_ptr[i]);
1090  }
1091 
1092  if (N % 2) {
1093 
1094  __half const *b_residual_ptr = reinterpret_cast<__half const *>(&b);
1095  __half const *c_residual_ptr = reinterpret_cast<__half const *>(&c);
1096  __half d_residual = __hfma(
1097  reinterpret_cast<__half const &>(a),
1098  b_residual_ptr[N - 1],
1099  c_residual_ptr[N - 1]);
1100 
1101  result[N - 1] = reinterpret_cast<half_t const &>(d_residual);
1102  }
1103 
1104  #else
1105 
1106  multiply_add<half_t> op;
1107 
1108  CUTLASS_PRAGMA_UNROLL
1109  for (int i = 0; i < N; ++i) {
1110  result[i] = op(a, b[i], c[i]);
1111  }
1112  #endif
1113 
1114  return result;
1115  }
1116 
1117  CUTLASS_HOST_DEVICE
1118  Array<half_t, N> operator()(
1119  Array<half_t, N> const &a,
1120  half_t const &b,
1121  Array<half_t, N> const &c) const {
1122 
1123  Array<half_t, N> result;
1124  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)
1125 
1126  __half2 *result_ptr = reinterpret_cast<__half2 *>(&result);
1127  __half2 const *a_ptr = reinterpret_cast<__half2 const *>(&a);
1128  __half2 b_pair = __half2half2(reinterpret_cast<__half const &>(b));
1129  __half2 const *c_ptr = reinterpret_cast<__half2 const *>(&c);
1130 
1131  CUTLASS_PRAGMA_UNROLL
1132  for (int i = 0; i < N / 2; ++i) {
1133  result_ptr[i] = __hfma2(a_ptr[i], b_pair, c_ptr[i]);
1134  }
1135 
1136  if (N % 2) {
1137 
1138  __half const *a_residual_ptr = reinterpret_cast<__half const *>(&a);
1139  __half const *c_residual_ptr = reinterpret_cast<__half const *>(&c);
1140 
1141  __half d_residual = __hfma(
1142  a_residual_ptr[N - 1],
1143  reinterpret_cast<__half const &>(b),
1144  c_residual_ptr[N - 1]);
1145 
1146  result[N - 1] = reinterpret_cast<half_t const &>(d_residual);
1147  }
1148 
1149  #else
1150 
1151  multiply_add<half_t> op;
1152 
1153  CUTLASS_PRAGMA_UNROLL
1154  for (int i = 0; i < N; ++i) {
1155  result[i] = op(a[i], b, c[i]);
1156  }
1157  #endif
1158 
1159  return result;
1160  }
1161 
1162  CUTLASS_HOST_DEVICE
1163  Array<half_t, N> operator()(
1164  Array<half_t, N> const &a,
1165  Array<half_t, N> const &b,
1166  half_t const &c) const {
1167 
1168  Array<half_t, N> result;
1169  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)
1170 
1171  __half2 *result_ptr = reinterpret_cast<__half2 *>(&result);
1172  __half2 const *a_ptr = reinterpret_cast<__half2 const *>(&a);
1173  __half2 const *b_ptr = reinterpret_cast<__half2 const *>(&b);
1174  __half2 c_pair = __half2half2(reinterpret_cast<__half const &>(c));
1175 
1176  CUTLASS_PRAGMA_UNROLL
1177  for (int i = 0; i < N / 2; ++i) {
1178  result_ptr[i] = __hfma2(a_ptr[i], b_ptr[i], c_pair);
1179  }
1180 
1181  if (N % 2) {
1182 
1183  __half const *a_residual_ptr = reinterpret_cast<__half const *>(&a);
1184  __half const *b_residual_ptr = reinterpret_cast<__half const *>(&b);
1185 
1186  __half d_residual = __hfma(
1187  a_residual_ptr[N - 1],
1188  b_residual_ptr[N - 1],
1189  reinterpret_cast<__half const &>(c));
1190 
1191  result[N - 1] = reinterpret_cast<half_t const &>(d_residual);
1192  }
1193 
1194  #else
1195 
1196  multiply_add<half_t> op;
1197 
1198  CUTLASS_PRAGMA_UNROLL
1199  for (int i = 0; i < N; ++i) {
1200  result[i] = op(a[i], b[i], c);
1201  }
1202  #endif
1203 
1204  return result;
1205  }
1206 };
1207 
1209 
1210 } // namespace cutlass
cutlass::multiply_add
Fused multiply-add.
Definition: functional.h:92
cutlass::minimum< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(T const &scalar, Array< T, N > const &rhs) const
Definition: functional.h:351
cutlass::plus< Array< half_t, N > >::operator()
CUTLASS_HOST_DEVICE Array< half_t, N > operator()(Array< half_t, N > const &lhs, Array< half_t, N > const &rhs) const
Definition: functional.h:578
cutlass
Definition: aligned_buffer.h:35
cutlass::xor_add::operator()
CUTLASS_HOST_DEVICE T operator()(T const &a, T const &b, T const &c) const
Definition: functional.h:103
cutlass::divides< Array< half_t, N > >::operator()
CUTLASS_HOST_DEVICE Array< half_t, N > operator()(Array< half_t, N > const &lhs, half_t const &rhs) const
Definition: functional.h:955
cutlass::minus< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(Array< T, N > const &lhs, T const &scalar) const
Definition: functional.h:383
cutlass::imag
CUTLASS_HOST_DEVICE float const & imag(cuFloatComplex const &z)
Returns the imaginary part of the complex number.
Definition: complex.h:72
cutlass::minus< Array< half_t, N > >::operator()
CUTLASS_HOST_DEVICE Array< half_t, N > operator()(Array< half_t, N > const &lhs, Array< half_t, N > const &rhs) const
Definition: functional.h:678
cutlass::plus::operator()
CUTLASS_HOST_DEVICE T operator()(T lhs, T const &rhs) const
Definition: functional.h:48
cutlass::maximum< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(Array< T, N > const &lhs, T const &scalar) const
Definition: functional.h:269
cutlass::minimum< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(Array< T, N > const &lhs, Array< T, N > const &rhs) const
Definition: functional.h:323
cutlass::plus< Array< half_t, N > >::operator()
CUTLASS_HOST_DEVICE Array< half_t, N > operator()(half_t const &lhs, Array< half_t, N > const &rhs) const
Definition: functional.h:611
half.h
Defines a class for using IEEE half-precision floating-point types in host or device code...
cutlass::divides< Array< half_t, N > >::operator()
CUTLASS_HOST_DEVICE Array< half_t, N > operator()(half_t const &lhs, Array< half_t, N > const &rhs) const
Definition: functional.h:920
cutlass::multiplies< Array< half_t, N > >::operator()
CUTLASS_HOST_DEVICE Array< half_t, N > operator()(Array< half_t, N > const &lhs, Array< half_t, N > const &rhs) const
Definition: functional.h:778
cutlass::plus< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(Array< T, N > const &lhs, Array< T, N > const &rhs) const
Definition: functional.h:191
cutlass::minimum
Definition: functional.h:298
cutlass::maximum< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(Array< T, N > const &lhs, Array< T, N > const &rhs) const
Definition: functional.h:255
cutlass::maximum
Definition: functional.h:235
cutlass::half_t
IEEE half-precision floating-point type.
Definition: half.h:126
cutlass::negate::operator()
CUTLASS_HOST_DEVICE T operator()(T lhs) const
Definition: functional.h:85
cutlass::multiplies< Array< half_t, N > >::operator()
CUTLASS_HOST_DEVICE Array< half_t, N > operator()(half_t const &lhs, Array< half_t, N > const &rhs) const
Definition: functional.h:811
cutlass::plus< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(Array< T, N > const &lhs, T const &scalar) const
Definition: functional.h:205
cutlass::real
CUTLASS_HOST_DEVICE float const & real(cuFloatComplex const &z)
Returns the real part of the complex number.
Definition: complex.h:56
cutlass::minus< Array< half_t, N > >::operator()
CUTLASS_HOST_DEVICE Array< half_t, N > operator()(Array< half_t, N > const &lhs, half_t const &rhs) const
Definition: functional.h:743
cutlass::minimum< float >::operator()
CUTLASS_HOST_DEVICE float operator()(float const &lhs, float const &rhs) const
Definition: functional.h:309
cutlass::minimum< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(Array< T, N > const &lhs, T const &scalar) const
Definition: functional.h:337
cutlass::complex::imag
CUTLASS_HOST_DEVICE T const & imag() const
Accesses the imaginary part of the complex number.
Definition: complex.h:240
cutlass::multiply_add::operator()
CUTLASS_HOST_DEVICE C operator()(A const &a, B const &b, C const &c) const
Definition: functional.h:94
cutlass::multiply_add< Array< half_t, N >, Array< half_t, N >, Array< half_t, N > >::operator()
CUTLASS_HOST_DEVICE Array< half_t, N > operator()(Array< half_t, N > const &a, Array< half_t, N > const &b, half_t const &c) const
Definition: functional.h:1163
cutlass::plus
Definition: functional.h:46
cutlass::minus< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(Array< T, N > const &lhs, Array< T, N > const &rhs) const
Definition: functional.h:369
cutlass::maximum::operator()
CUTLASS_HOST_DEVICE T operator()(T const &lhs, T const &rhs) const
Definition: functional.h:238
cutlass::multiplies< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(Array< T, N > const &lhs, Array< T, N > const &rhs) const
Definition: functional.h:415
cutlass::minus< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(T const &scalar, Array< T, N > const &rhs) const
Definition: functional.h:397
cutlass::multiplies::operator()
CUTLASS_HOST_DEVICE T operator()(T lhs, T const &rhs) const
Definition: functional.h:66
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::maximum< float >::operator()
CUTLASS_HOST_DEVICE float operator()(float const &lhs, float const &rhs) const
Definition: functional.h:246
cutlass::negate
Definition: functional.h:83
cutlass::divides::operator()
CUTLASS_HOST_DEVICE T operator()(T lhs, T const &rhs) const
Definition: functional.h:75
cutlass::multiply_add< Array< T, N >, Array< T, N >, Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(Array< T, N > const &a, T const &scalar, Array< T, N > const &c) const
Definition: functional.h:541
cutlass::multiply_add< Array< T, N >, Array< T, N >, Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(Array< T, N > const &a, Array< T, N > const &b, Array< T, N > const &c) const
Definition: functional.h:527
cutlass::multiplies< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(T const &scalar, Array< T, N > const &rhs) const
Definition: functional.h:443
cutlass::multiplies
Definition: functional.h:64
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:89
numeric_types.h
Top-level include for all CUTLASS numeric types.
cutlass::multiplies< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(Array< T, N > const &lhs, T const &scalar) const
Definition: functional.h:429
cutlass::divides< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(T const &scalar, Array< T, N > const &rhs) const
Definition: functional.h:489
cutlass::divides< Array< half_t, N > >::operator()
CUTLASS_HOST_DEVICE Array< half_t, N > operator()(Array< half_t, N > const &lhs, Array< half_t, N > const &rhs) const
Definition: functional.h:884
cutlass::divides
Definition: functional.h:73
cutlass::complex::real
CUTLASS_HOST_DEVICE T const & real() const
Accesses the real part of the complex number.
Definition: complex.h:232
cutlass::plus< Array< half_t, N > >::operator()
CUTLASS_HOST_DEVICE Array< half_t, N > operator()(Array< half_t, N > const &lhs, half_t const &rhs) const
Definition: functional.h:643
cutlass::complex
Definition: complex.h:92
cutlass::minimum::operator()
CUTLASS_HOST_DEVICE T operator()(T const &lhs, T const &rhs) const
Definition: functional.h:301
cutlass::minimum< Array< T, N > >::scalar_op
static CUTLASS_HOST_DEVICE T scalar_op(T const &lhs, T const &rhs)
Definition: functional.h:318
cutlass::multiply_add< T, complex< T >, complex< T > >::operator()
CUTLASS_HOST_DEVICE complex< T > operator()(T const &a, complex< T > const &b, complex< T > const &c) const
Definition: functional.h:164
cutlass::multiply_add< Array< half_t, N >, Array< half_t, N >, Array< half_t, N > >::operator()
CUTLASS_HOST_DEVICE Array< half_t, N > operator()(half_t const &a, Array< half_t, N > const &b, Array< half_t, N > const &c) const
Definition: functional.h:1074
cutlass::multiply_add< complex< T >, T, complex< T > >::operator()
CUTLASS_HOST_DEVICE complex< T > operator()(complex< T > const &a, T const &b, complex< T > const &c) const
Definition: functional.h:142
cutlass::negate< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(Array< T, N > const &lhs) const
Definition: functional.h:508
cutlass::xor_add
Fused multiply-add.
Definition: functional.h:101
cutlass::maximum< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(T const &scalar, Array< T, N > const &rhs) const
Definition: functional.h:283
cutlass::divides< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(Array< T, N > const &lhs, T const &scalar) const
Definition: functional.h:475
cutlass::minus::operator()
CUTLASS_HOST_DEVICE T operator()(T lhs, T const &rhs) const
Definition: functional.h:57
cutlass::multiply_add< Array< half_t, N >, Array< half_t, N >, Array< half_t, N > >::operator()
CUTLASS_HOST_DEVICE Array< half_t, N > operator()(Array< half_t, N > const &a, half_t const &b, Array< half_t, N > const &c) const
Definition: functional.h:1118
cutlass::plus< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(T const &scalar, Array< T, N > const &rhs) const
Definition: functional.h:219
cutlass::minus< Array< half_t, N > >::operator()
CUTLASS_HOST_DEVICE Array< half_t, N > operator()(half_t const &lhs, Array< half_t, N > const &rhs) const
Definition: functional.h:711
cutlass::multiply_add< Array< T, N >, Array< T, N >, Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(T const &scalar, Array< T, N > const &b, Array< T, N > const &c) const
Definition: functional.h:555
cutlass::divides< Array< T, N > >::operator()
CUTLASS_HOST_DEVICE Array< T, N > operator()(Array< T, N > const &lhs, Array< T, N > const &rhs) const
Definition: functional.h:461
cutlass::negate< Array< half_t, N > >::operator()
CUTLASS_HOST_DEVICE Array< half_t, N > operator()(Array< half_t, N > const &lhs) const
Definition: functional.h:993
cutlass::minus
Definition: functional.h:55
cutlass::multiply_add< complex< T >, complex< T >, complex< T > >::operator()
CUTLASS_HOST_DEVICE complex< T > operator()(complex< T > const &a, complex< T > const &b, complex< T > const &c) const
Definition: functional.h:118
cutlass::multiplies< Array< half_t, N > >::operator()
CUTLASS_HOST_DEVICE Array< half_t, N > operator()(Array< half_t, N > const &lhs, half_t const &rhs) const
Definition: functional.h:846
cutlass.h
Basic include for CUTLASS.
cutlass::multiply_add< Array< half_t, N >, Array< half_t, N >, Array< half_t, N > >::operator()
CUTLASS_HOST_DEVICE Array< half_t, N > operator()(Array< half_t, N > const &a, Array< half_t, N > const &b, Array< half_t, N > const &c) const
Definition: functional.h:1028
Generated by   doxygen 1.8.11