/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/library/utility/host_tensor_generator.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/library/utility/host_tensor_generator.hpp Source File#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/library/utility/host_tensor_generator.hpp Source File
host_tensor_generator.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 
6 #include <cmath>
7 #include <numeric>
8 #include <random>
9 
10 #include "ck/ck.hpp"
11 
12 template <typename T>
14 {
15  template <typename... Is>
16  T operator()(Is...)
17  {
18  return T{0};
19  }
20 };
21 
22 template <typename T>
24 {
25  T value = 1;
26 
27  template <typename... Is>
28  T operator()(Is...)
29  {
30  return value;
31  }
32 };
33 
34 template <>
36 {
37  float value = 1.0;
38 
39  template <typename... Is>
41  {
42  return ck::type_convert<ck::half_t>(value);
43  }
44 };
45 
46 template <>
48 {
49  float value = 1.0;
50 
51  template <typename... Is>
53  {
54  return ck::type_convert<ck::bhalf_t>(value);
55  }
56 };
57 
58 #if defined CK_ENABLE_FP8
59 template <>
60 struct GeneratorTensor_1<ck::f8_t>
61 {
62  float value = 1.0;
63 
64  template <typename... Is>
65  ck::f8_t operator()(Is...)
66  {
67  return ck::type_convert<ck::f8_t>(value);
68  }
69 };
70 
71 template <>
72 struct GeneratorTensor_1<ck::bf8_t>
73 {
74  float value = 1.0;
75 
76  template <typename... Is>
77  ck::bf8_t operator()(Is...)
78  {
79  return ck::type_convert<ck::bf8_t>(value);
80  }
81 };
82 #endif
83 
84 template <>
86 {
87  float value = 1.0;
88 
89  template <typename... Is>
91  {
92  return ck::type_convert<ck::f4_t>(value);
93  }
94 };
95 
96 template <>
97 struct GeneratorTensor_1<ck::f4x2_pk_t>
98 {
99  float value = 1.0;
100 
101  template <typename... Is>
103  {
104  return ck::f4x2_pk_t{ck::type_convert<ck::f4x2_t>(ck::float2_t{value, value})};
105  }
106 };
107 
108 template <>
110 {
111  float value = 1.0;
112 
113  template <typename... Is>
115  {
116  ck::f6x32_pk_t r;
117  ck::static_for<0, 32, 1>{}([&](auto i) {
118  r.pack(ck::type_convert<ck::f6_t>(value), static_cast<ck::index_t>(i));
119  });
120  return r;
121  }
122 };
123 
124 template <>
126 {
127  float value = 1.0;
128 
129  template <typename... Is>
131  {
132  ck::bf6x32_pk_t r;
133  ck::static_for<0, 32, 1>{}([&](auto i) {
134  r.pack(ck::type_convert<ck::bf6_t>(value), static_cast<ck::index_t>(i));
135  });
136  return r;
137  }
138 };
139 
140 template <>
142 {
144 
145  template <typename... Is>
147  {
148  return value;
149  }
150 };
151 
152 template <>
153 struct GeneratorTensor_1<ck::pk_i4_t>
154 {
156 
157  template <typename... Is>
159  {
160  int t = value + 8;
161  ck::pk_i4_t r = ((t << 4) + t) & 0xff;
162  return r;
163  }
164 };
165 
166 template <>
167 struct GeneratorTensor_1<ck::e8m0_bexp_t>
168 {
169  float value = 1;
170 
171  template <typename... Is>
173  {
174  return ck::type_convert<ck::e8m0_bexp_t>(value);
175  }
176 };
177 
178 template <typename T>
180 {
181  int min_value = 0;
182  int max_value = 1;
183 
184  template <typename... Is>
185  T operator()(Is...)
186  {
187  return static_cast<T>((std::rand() % (max_value - min_value)) + min_value);
188  }
189 };
190 
191 template <>
193 {
194  int min_value = 0;
195  int max_value = 1;
196 
197  template <typename... Is>
199  {
200  ck::f6x32_pk_t r;
201  ck::static_for<0, 32, 1>{}([&](auto i) {
202  float tmp = (std::rand() % (max_value - min_value)) + min_value;
203  r.pack(ck::type_convert<ck::f6_t>(tmp), static_cast<ck::index_t>(i));
204  });
205 
206  return r;
207  }
208 };
209 
210 template <>
212 {
213  int min_value = 0;
214  int max_value = 1;
215 
216  template <typename... Is>
218  {
219  ck::bf6x32_pk_t r;
220  ck::static_for<0, 32, 1>{}([&](auto i) {
221  float tmp = (std::rand() % (max_value - min_value)) + min_value;
222  r.pack(ck::type_convert<ck::bf6_t>(tmp), static_cast<ck::index_t>(i));
223  });
224 
225  return r;
226  }
227 };
228 
229 template <>
231 {
232  int min_value = 0;
233  int max_value = 1;
234 
235  template <typename... Is>
237  {
238  float tmp = (std::rand() % (max_value - min_value)) + min_value;
239  return ck::type_convert<ck::bhalf_t>(tmp);
240  }
241 };
242 
243 template <>
245 {
246  int min_value = 0;
247  int max_value = 1;
248 
249  template <typename... Is>
251  {
252  return (std::rand() % (max_value - min_value)) + min_value;
253  }
254 };
255 
256 template <>
257 struct GeneratorTensor_2<ck::pk_i4_t>
258 {
259  int min_value = 0;
260  int max_value = 1;
261 
262  template <typename... Is>
264  {
265  int hi = std::rand() % (max_value - min_value) + min_value + 8;
266  int lo = std::rand() % (max_value - min_value) + min_value + 8;
267  ck::pk_i4_t r = (((hi & 0xf) << 4) + (lo & 0xf));
268  return r;
269  }
270 };
271 
272 #if defined CK_ENABLE_FP8
273 template <>
274 struct GeneratorTensor_2<ck::f8_t>
275 {
276  int min_value = 0;
277  int max_value = 1;
278 
279  template <typename... Is>
280  ck::f8_t operator()(Is...)
281  {
282  float tmp = (std::rand() % (max_value - min_value)) + min_value;
283  return ck::type_convert<ck::f8_t>(tmp);
284  }
285 };
286 #endif
287 
288 #if defined CK_ENABLE_BF8
289 template <>
290 struct GeneratorTensor_2<ck::bf8_t>
291 {
292  int min_value = 0;
293  int max_value = 1;
294 
295  template <typename... Is>
296  ck::bf8_t operator()(Is...)
297  {
298  float tmp = (std::rand() % (max_value - min_value)) + min_value;
299  return ck::type_convert<ck::bf8_t>(tmp);
300  }
301 };
302 #endif
303 
304 template <>
306 {
307  int min_value = 0;
308  int max_value = 1;
309 
310  template <typename... Is>
312  {
313  float tmp = (std::rand() % (max_value - min_value)) + min_value;
314  return ck::type_convert<ck::f4_t>(tmp);
315  }
316 };
317 
318 template <>
319 struct GeneratorTensor_2<ck::f4x2_pk_t>
320 {
321  int min_value = 0;
322  int max_value = 1;
323 
324  template <typename... Is>
326  {
327  float tmp0 = (std::rand() % (max_value - min_value)) + min_value;
328  float tmp1 = (std::rand() % (max_value - min_value)) + min_value;
329  return ck::f4x2_pk_t{ck::type_convert<ck::f4x2_t>(ck::float2_t{tmp0, tmp1})};
330  }
331 };
332 
333 template <typename T>
335 {
336  float min_value = 0;
337  float max_value = 1;
338 
339  template <typename... Is>
340  T operator()(Is...)
341  {
342  float tmp = float(std::rand()) / float(RAND_MAX);
343 
344  return static_cast<T>(min_value + tmp * (max_value - min_value));
345  }
346 };
347 
348 template <>
350 {
351  float min_value = 0;
352  float max_value = 1;
353 
354  template <typename... Is>
356  {
357  float tmp = float(std::rand()) / float(RAND_MAX);
358 
359  float fp32_tmp = min_value + tmp * (max_value - min_value);
360 
361  return ck::type_convert<ck::bhalf_t>(fp32_tmp);
362  }
363 };
364 
365 #if defined CK_ENABLE_FP8
366 template <>
367 struct GeneratorTensor_3<ck::f8_t>
368 {
369  float min_value = 0;
370  float max_value = 1;
371 
372  template <typename... Is>
373  ck::f8_t operator()(Is...)
374  {
375  float tmp = float(std::rand()) / float(RAND_MAX);
376 
377  float fp32_tmp = min_value + tmp * (max_value - min_value);
378 
379  return ck::type_convert<ck::f8_t>(fp32_tmp);
380  }
381 };
382 #endif
383 
384 #if defined CK_ENABLE_BF8
385 template <>
386 struct GeneratorTensor_3<ck::bf8_t>
387 {
388  float min_value = 0;
389  float max_value = 1;
390 
391  template <typename... Is>
392  ck::bf8_t operator()(Is...)
393  {
394  float tmp = float(std::rand()) / float(RAND_MAX);
395 
396  float fp32_tmp = min_value + tmp * (max_value - min_value);
397 
398  return ck::type_convert<ck::bf8_t>(fp32_tmp);
399  }
400 };
401 #endif
402 
403 template <>
405 {
406  float min_value = 0;
407  float max_value = 1;
408 
409  template <typename... Is>
411  {
412  float tmp = float(std::rand()) / float(RAND_MAX);
413 
414  float fp32_tmp = min_value + tmp * (max_value - min_value);
415 
416  return ck::type_convert<ck::f4_t>(fp32_tmp);
417  }
418 };
419 
420 template <>
421 struct GeneratorTensor_3<ck::f4x2_pk_t>
422 {
423  float min_value = 0;
424  float max_value = 1;
425 
426  template <typename... Is>
428  {
429  float tmp0 = float(std::rand()) / float(RAND_MAX);
430  float tmp1 = float(std::rand()) / float(RAND_MAX);
431 
432  float fp32_tmp0 = min_value + tmp0 * (max_value - min_value);
433  float fp32_tmp1 = min_value + tmp1 * (max_value - min_value);
434 
435  return ck::f4x2_pk_t{ck::type_convert<ck::f4x2_t>(ck::float2_t{fp32_tmp0, fp32_tmp1})};
436  }
437 };
438 
439 template <>
440 struct GeneratorTensor_3<ck::pk_i4_t>
441 {
442  int min_value = 0;
443  int max_value = 1;
444 
445  template <typename... Is>
447  {
448  int hi = std::rand() % (max_value - min_value) + min_value + 8;
449  int lo = std::rand() % (max_value - min_value) + min_value + 8;
450  ck::pk_i4_t r = (((hi & 0xf) << 4) + (lo & 0xf));
451  return r;
452  }
453 };
454 
455 template <>
457 {
458  float min_value = 0;
459  float max_value = 1;
460 
461  template <typename... Is>
463  {
464  ck::f6x32_pk_t r;
465  ck::static_for<0, 32, 1>{}([&](auto i) {
466  float rnd = float(std::rand()) / float(RAND_MAX);
467  float fp32 = min_value + rnd * (max_value - min_value);
468  r.pack(ck::type_convert<ck::f6_t>(fp32), static_cast<ck::index_t>(i));
469  });
470 
471  return r;
472  }
473 };
474 
475 template <>
477 {
478  float min_value = 0;
479  float max_value = 1;
480 
481  template <typename... Is>
483  {
484  ck::bf6x32_pk_t r;
485  ck::static_for<0, 32, 1>{}([&](auto i) {
486  float rnd = float(std::rand()) / float(RAND_MAX);
487  float fp32 = min_value + rnd * (max_value - min_value);
488  r.pack(ck::type_convert<ck::bf6_t>(fp32), static_cast<ck::index_t>(i));
489  });
490 
491  return r;
492  }
493 };
494 
495 template <typename T>
497 {
498  std::mt19937 generator;
499  std::normal_distribution<float> distribution;
500 
501  GeneratorTensor_4(float mean, float stddev, unsigned int seed = 1)
502  : generator(seed), distribution(mean, stddev){};
503 
504  template <typename... Is>
505  T operator()(Is...)
506  {
507  float tmp = distribution(generator);
508 
509  return ck::type_convert<T>(tmp);
510  }
511 };
512 
513 template <>
514 struct GeneratorTensor_4<ck::f4x2_pk_t>
515 {
516  std::mt19937 generator;
517  std::normal_distribution<float> distribution;
518 
519  GeneratorTensor_4(float mean, float stddev, unsigned int seed = 1)
520  : generator(seed), distribution(mean, stddev){};
521 
522  template <typename... Is>
524  {
525  float fp32_tmp0 = distribution(generator);
526  float fp32_tmp1 = distribution(generator);
527 
528  return ck::f4x2_pk_t{ck::type_convert<ck::f4x2_t>(ck::float2_t{fp32_tmp0, fp32_tmp1})};
529  }
530 };
531 
532 template <>
534 {
535  std::mt19937 generator;
536  std::normal_distribution<float> distribution;
537 
538  GeneratorTensor_4(float mean, float stddev, unsigned int seed = 1)
539  : generator(seed), distribution(mean, stddev){};
540 
541  template <typename... Is>
543  {
544  ck::f6x32_pk_t r;
545  ck::static_for<0, 32, 1>{}([&](auto i) {
546  r.pack(ck::type_convert<ck::f6_t>(distribution(generator)),
547  static_cast<ck::index_t>(i));
548  });
549 
550  return r;
551  }
552 };
553 
554 template <>
556 {
557  std::mt19937 generator;
558  std::normal_distribution<float> distribution;
559 
560  GeneratorTensor_4(float mean, float stddev, unsigned int seed = 1)
561  : generator(seed), distribution(mean, stddev){};
562 
563  template <typename... Is>
565  {
566  ck::bf6x32_pk_t r;
567  ck::static_for<0, 32, 1>{}([&](auto i) {
568  r.pack(ck::type_convert<ck::bf6_t>(distribution(generator)),
569  static_cast<ck::index_t>(i));
570  });
571 
572  return r;
573  }
574 };
575 
577 {
578  template <typename... Ts>
579  float operator()(Ts... Xs) const
580  {
581  std::array<ck::index_t, sizeof...(Ts)> dims = {static_cast<ck::index_t>(Xs)...};
582  return std::accumulate(dims.begin(),
583  dims.end(),
584  true,
585  [](bool init, ck::index_t x) -> int { return init != (x % 2); })
586  ? 1
587  : -1;
588  }
589 };
590 
608 template <typename T, ck::index_t Dim>
610 {
611  template <typename... Ts>
612  T operator()(Ts... Xs) const
613  {
614  std::array<ck::index_t, sizeof...(Ts)> dims = {{static_cast<ck::index_t>(Xs)...}};
615 
616  float tmp = dims[Dim];
617  return ck::type_convert<T>(tmp);
618  }
619 };
620 
621 template <ck::index_t Dim>
622 struct GeneratorTensor_Sequential<ck::f4x2_pk_t, Dim>
623 {
624  template <typename... Ts>
625  ck::f4x2_pk_t operator()(Ts... Xs) const
626  {
627  std::array<ck::index_t, sizeof...(Ts)> dims = {{static_cast<ck::index_t>(Xs)...}};
628 
629  float tmp = dims[Dim];
630  return ck::type_convert<ck::f4x2_t>(ck::float2_t(tmp));
631  }
632 };
633 
634 template <ck::index_t Dim>
636 {
637  template <typename... Ts>
638  ck::f6x32_pk_t operator()(Ts... Xs) const
639  {
640  std::array<ck::index_t, sizeof...(Ts)> dims = {{static_cast<ck::index_t>(Xs)...}};
641 
642  float tmp = dims[Dim];
643 
644  ck::f6x32_pk_t r;
646  [&](auto i) { r.pack(ck::type_convert<ck::f6_t>(tmp), static_cast<ck::index_t>(i)); });
647  return r;
648  }
649 };
650 
651 template <ck::index_t Dim>
653 {
654  template <typename... Ts>
655  ck::bf6x32_pk_t operator()(Ts... Xs) const
656  {
657  std::array<ck::index_t, sizeof...(Ts)> dims = {{static_cast<ck::index_t>(Xs)...}};
658 
659  float tmp = dims[Dim];
660 
661  ck::bf6x32_pk_t r;
663  [&](auto i) { r.pack(ck::type_convert<ck::bf6_t>(tmp), static_cast<ck::index_t>(i)); });
664  return r;
665  }
666 };
667 
668 template <typename T, size_t NumEffectiveDim = 2>
670 {
671  T value{1};
672 
673  template <typename... Ts>
674  T operator()(Ts... Xs) const
675  {
676  std::array<ck::index_t, sizeof...(Ts)> dims = {{static_cast<ck::index_t>(Xs)...}};
677  size_t start_dim = dims.size() - NumEffectiveDim;
678  bool pred = true;
679  for(size_t i = start_dim + 1; i < dims.size(); i++)
680  {
681  pred &= (dims[start_dim] == dims[i]);
682  }
683  return pred ? value : T{0};
684  }
685 };
Definition: ck.hpp:268
bf8_fnuz_t bf8_t
Definition: amd_ck_fp8.hpp:1763
f8_fnuz_t f8_t
Definition: amd_ck_fp8.hpp:1762
unsigned _BitInt(4) f4_t
Definition: data_type.hpp:33
typename vector_type< float, 2 >::type float2_t
Definition: dtype_vector.hpp:2145
_Float16 half_t
Definition: data_type.hpp:31
ushort bhalf_t
Definition: data_type.hpp:30
f6_pk_t< f6_t, 32 > f6x32_pk_t
Definition: data_type.hpp:181
int32_t index_t
Definition: ck.hpp:299
f6_pk_t< bf6_t, 32 > bf6x32_pk_t
Definition: data_type.hpp:183
signed char int8_t
Definition: stdint.h:121
Definition: host_tensor_generator.hpp:14
T operator()(Is...)
Definition: host_tensor_generator.hpp:16
ck::bf6x32_pk_t operator()(Is...)
Definition: host_tensor_generator.hpp:130
ck::bhalf_t operator()(Is...)
Definition: host_tensor_generator.hpp:52
ck::e8m0_bexp_t operator()(Is...)
Definition: host_tensor_generator.hpp:172
ck::f4_t operator()(Is...)
Definition: host_tensor_generator.hpp:90
ck::f4x2_pk_t operator()(Is...)
Definition: host_tensor_generator.hpp:102
ck::f6x32_pk_t operator()(Is...)
Definition: host_tensor_generator.hpp:114
ck::half_t operator()(Is...)
Definition: host_tensor_generator.hpp:40
ck::pk_i4_t operator()(Is...)
Definition: host_tensor_generator.hpp:158
int8_t operator()(Is...)
Definition: host_tensor_generator.hpp:146
Definition: host_tensor_generator.hpp:24
T value
Definition: host_tensor_generator.hpp:25
T operator()(Is...)
Definition: host_tensor_generator.hpp:28
ck::bf6x32_pk_t operator()(Is...)
Definition: host_tensor_generator.hpp:217
ck::bhalf_t operator()(Is...)
Definition: host_tensor_generator.hpp:236
ck::f4_t operator()(Is...)
Definition: host_tensor_generator.hpp:311
ck::f4x2_pk_t operator()(Is...)
Definition: host_tensor_generator.hpp:325
ck::f6x32_pk_t operator()(Is...)
Definition: host_tensor_generator.hpp:198
ck::pk_i4_t operator()(Is...)
Definition: host_tensor_generator.hpp:263
int8_t operator()(Is...)
Definition: host_tensor_generator.hpp:250
Definition: host_tensor_generator.hpp:180
int max_value
Definition: host_tensor_generator.hpp:182
int min_value
Definition: host_tensor_generator.hpp:181
T operator()(Is...)
Definition: host_tensor_generator.hpp:185
ck::bf6x32_pk_t operator()(Is...)
Definition: host_tensor_generator.hpp:482
ck::bhalf_t operator()(Is...)
Definition: host_tensor_generator.hpp:355
ck::f4_t operator()(Is...)
Definition: host_tensor_generator.hpp:410
ck::f4x2_pk_t operator()(Is...)
Definition: host_tensor_generator.hpp:427
ck::f6x32_pk_t operator()(Is...)
Definition: host_tensor_generator.hpp:462
ck::pk_i4_t operator()(Is...)
Definition: host_tensor_generator.hpp:446
Definition: host_tensor_generator.hpp:335
float max_value
Definition: host_tensor_generator.hpp:337
float min_value
Definition: host_tensor_generator.hpp:336
T operator()(Is...)
Definition: host_tensor_generator.hpp:340
std::normal_distribution< float > distribution
Definition: host_tensor_generator.hpp:558
ck::bf6x32_pk_t operator()(Is...)
Definition: host_tensor_generator.hpp:564
std::mt19937 generator
Definition: host_tensor_generator.hpp:557
GeneratorTensor_4(float mean, float stddev, unsigned int seed=1)
Definition: host_tensor_generator.hpp:560
GeneratorTensor_4(float mean, float stddev, unsigned int seed=1)
Definition: host_tensor_generator.hpp:519
ck::f4x2_pk_t operator()(Is...)
Definition: host_tensor_generator.hpp:523
std::normal_distribution< float > distribution
Definition: host_tensor_generator.hpp:517
std::mt19937 generator
Definition: host_tensor_generator.hpp:516
GeneratorTensor_4(float mean, float stddev, unsigned int seed=1)
Definition: host_tensor_generator.hpp:538
std::mt19937 generator
Definition: host_tensor_generator.hpp:535
std::normal_distribution< float > distribution
Definition: host_tensor_generator.hpp:536
ck::f6x32_pk_t operator()(Is...)
Definition: host_tensor_generator.hpp:542
Definition: host_tensor_generator.hpp:497
std::mt19937 generator
Definition: host_tensor_generator.hpp:498
GeneratorTensor_4(float mean, float stddev, unsigned int seed=1)
Definition: host_tensor_generator.hpp:501
T operator()(Is...)
Definition: host_tensor_generator.hpp:505
std::normal_distribution< float > distribution
Definition: host_tensor_generator.hpp:499
Definition: host_tensor_generator.hpp:577
float operator()(Ts... Xs) const
Definition: host_tensor_generator.hpp:579
Definition: host_tensor_generator.hpp:670
T operator()(Ts... Xs) const
Definition: host_tensor_generator.hpp:674
T value
Definition: host_tensor_generator.hpp:671
ck::bf6x32_pk_t operator()(Ts... Xs) const
Definition: host_tensor_generator.hpp:655
ck::f4x2_pk_t operator()(Ts... Xs) const
Definition: host_tensor_generator.hpp:625
ck::f6x32_pk_t operator()(Ts... Xs) const
Definition: host_tensor_generator.hpp:638
Is used to generate sequential values based on the specified dimension.
Definition: host_tensor_generator.hpp:610
T operator()(Ts... Xs) const
Definition: host_tensor_generator.hpp:612
Definition: amd_ck_fp8.hpp:49
Unsigned representation of a conventional biased Float32 exponent.
Definition: e8m0.hpp:26
Definition: data_type.hpp:42
Definition: data_type.hpp:79
__host__ __device__ void pack(const T x, const index_t i)
Definition: data_type.hpp:119
Definition: amd_ck_fp8.hpp:36
Definition: data_type.hpp:187
Definition: functional2.hpp:33