/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 << 4) + lo) & 0xff;
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 <>
441 {
442  float min_value = 0;
443  float max_value = 1;
444 
445  template <typename... Is>
447  {
448  ck::f6x32_pk_t r;
449  ck::static_for<0, 32, 1>{}([&](auto i) {
450  float rnd = float(std::rand()) / float(RAND_MAX);
451  float fp32 = min_value + rnd * (max_value - min_value);
452  r.pack(ck::type_convert<ck::f6_t>(fp32), static_cast<ck::index_t>(i));
453  });
454 
455  return r;
456  }
457 };
458 
459 template <>
461 {
462  float min_value = 0;
463  float max_value = 1;
464 
465  template <typename... Is>
467  {
468  ck::bf6x32_pk_t r;
469  ck::static_for<0, 32, 1>{}([&](auto i) {
470  float rnd = float(std::rand()) / float(RAND_MAX);
471  float fp32 = min_value + rnd * (max_value - min_value);
472  r.pack(ck::type_convert<ck::bf6_t>(fp32), static_cast<ck::index_t>(i));
473  });
474 
475  return r;
476  }
477 };
478 
479 template <typename T>
481 {
482  std::mt19937 generator;
483  std::normal_distribution<float> distribution;
484 
485  GeneratorTensor_4(float mean, float stddev, unsigned int seed = 1)
486  : generator(seed), distribution(mean, stddev){};
487 
488  template <typename... Is>
489  T operator()(Is...)
490  {
491  float tmp = distribution(generator);
492 
493  return ck::type_convert<T>(tmp);
494  }
495 };
496 
497 template <>
498 struct GeneratorTensor_4<ck::f4x2_pk_t>
499 {
500  std::mt19937 generator;
501  std::normal_distribution<float> distribution;
502 
503  GeneratorTensor_4(float mean, float stddev, unsigned int seed = 1)
504  : generator(seed), distribution(mean, stddev){};
505 
506  template <typename... Is>
508  {
509  float fp32_tmp0 = distribution(generator);
510  float fp32_tmp1 = distribution(generator);
511 
512  return ck::f4x2_pk_t{ck::type_convert<ck::f4x2_t>(ck::float2_t{fp32_tmp0, fp32_tmp1})};
513  }
514 };
515 
516 template <>
518 {
519  std::mt19937 generator;
520  std::normal_distribution<float> distribution;
521 
522  GeneratorTensor_4(float mean, float stddev, unsigned int seed = 1)
523  : generator(seed), distribution(mean, stddev){};
524 
525  template <typename... Is>
527  {
528  ck::f6x32_pk_t r;
529  ck::static_for<0, 32, 1>{}([&](auto i) {
530  r.pack(ck::type_convert<ck::f6_t>(distribution(generator)),
531  static_cast<ck::index_t>(i));
532  });
533 
534  return r;
535  }
536 };
537 
538 template <>
540 {
541  std::mt19937 generator;
542  std::normal_distribution<float> distribution;
543 
544  GeneratorTensor_4(float mean, float stddev, unsigned int seed = 1)
545  : generator(seed), distribution(mean, stddev){};
546 
547  template <typename... Is>
549  {
550  ck::bf6x32_pk_t r;
551  ck::static_for<0, 32, 1>{}([&](auto i) {
552  r.pack(ck::type_convert<ck::bf6_t>(distribution(generator)),
553  static_cast<ck::index_t>(i));
554  });
555 
556  return r;
557  }
558 };
559 
561 {
562  template <typename... Ts>
563  float operator()(Ts... Xs) const
564  {
565  std::array<ck::index_t, sizeof...(Ts)> dims = {static_cast<ck::index_t>(Xs)...};
566  return std::accumulate(dims.begin(),
567  dims.end(),
568  true,
569  [](bool init, ck::index_t x) -> int { return init != (x % 2); })
570  ? 1
571  : -1;
572  }
573 };
574 
592 template <typename T, ck::index_t Dim>
594 {
595  template <typename... Ts>
596  T operator()(Ts... Xs) const
597  {
598  std::array<ck::index_t, sizeof...(Ts)> dims = {{static_cast<ck::index_t>(Xs)...}};
599 
600  float tmp = dims[Dim];
601  return ck::type_convert<T>(tmp);
602  }
603 };
604 
605 template <ck::index_t Dim>
606 struct GeneratorTensor_Sequential<ck::f4x2_pk_t, Dim>
607 {
608  template <typename... Ts>
609  ck::f4x2_pk_t operator()(Ts... Xs) const
610  {
611  std::array<ck::index_t, sizeof...(Ts)> dims = {{static_cast<ck::index_t>(Xs)...}};
612 
613  float tmp = dims[Dim];
614  return ck::type_convert<ck::f4x2_t>(ck::float2_t(tmp));
615  }
616 };
617 
618 template <ck::index_t Dim>
620 {
621  template <typename... Ts>
622  ck::f6x32_pk_t operator()(Ts... Xs) const
623  {
624  std::array<ck::index_t, sizeof...(Ts)> dims = {{static_cast<ck::index_t>(Xs)...}};
625 
626  float tmp = dims[Dim];
627 
628  ck::f6x32_pk_t r;
630  [&](auto i) { r.pack(ck::type_convert<ck::f6_t>(tmp), static_cast<ck::index_t>(i)); });
631  return r;
632  }
633 };
634 
635 template <ck::index_t Dim>
637 {
638  template <typename... Ts>
639  ck::bf6x32_pk_t operator()(Ts... Xs) const
640  {
641  std::array<ck::index_t, sizeof...(Ts)> dims = {{static_cast<ck::index_t>(Xs)...}};
642 
643  float tmp = dims[Dim];
644 
645  ck::bf6x32_pk_t r;
647  [&](auto i) { r.pack(ck::type_convert<ck::bf6_t>(tmp), static_cast<ck::index_t>(i)); });
648  return r;
649  }
650 };
651 
652 template <typename T, size_t NumEffectiveDim = 2>
654 {
655  T value{1};
656 
657  template <typename... Ts>
658  T operator()(Ts... Xs) const
659  {
660  std::array<ck::index_t, sizeof...(Ts)> dims = {{static_cast<ck::index_t>(Xs)...}};
661  size_t start_dim = dims.size() - NumEffectiveDim;
662  bool pred = true;
663  for(size_t i = start_dim + 1; i < dims.size(); i++)
664  {
665  pred &= (dims[start_dim] == dims[i]);
666  }
667  return pred ? value : T{0};
668  }
669 };
Definition: ck.hpp:267
bf8_fnuz_t bf8_t
Definition: amd_ck_fp8.hpp:1738
f8_fnuz_t f8_t
Definition: amd_ck_fp8.hpp:1737
unsigned _BitInt(4) f4_t
Definition: data_type.hpp:32
typename vector_type< float, 2 >::type float2_t
Definition: dtype_vector.hpp:2131
_Float16 half_t
Definition: data_type.hpp:30
ushort bhalf_t
Definition: data_type.hpp:29
f6_pk_t< f6_t, 32 > f6x32_pk_t
Definition: data_type.hpp:180
int32_t index_t
Definition: ck.hpp:298
f6_pk_t< bf6_t, 32 > bf6x32_pk_t
Definition: data_type.hpp:182
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:466
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: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:542
ck::bf6x32_pk_t operator()(Is...)
Definition: host_tensor_generator.hpp:548
std::mt19937 generator
Definition: host_tensor_generator.hpp:541
GeneratorTensor_4(float mean, float stddev, unsigned int seed=1)
Definition: host_tensor_generator.hpp:544
GeneratorTensor_4(float mean, float stddev, unsigned int seed=1)
Definition: host_tensor_generator.hpp:503
ck::f4x2_pk_t operator()(Is...)
Definition: host_tensor_generator.hpp:507
std::normal_distribution< float > distribution
Definition: host_tensor_generator.hpp:501
std::mt19937 generator
Definition: host_tensor_generator.hpp:500
GeneratorTensor_4(float mean, float stddev, unsigned int seed=1)
Definition: host_tensor_generator.hpp:522
std::mt19937 generator
Definition: host_tensor_generator.hpp:519
std::normal_distribution< float > distribution
Definition: host_tensor_generator.hpp:520
ck::f6x32_pk_t operator()(Is...)
Definition: host_tensor_generator.hpp:526
Definition: host_tensor_generator.hpp:481
std::mt19937 generator
Definition: host_tensor_generator.hpp:482
GeneratorTensor_4(float mean, float stddev, unsigned int seed=1)
Definition: host_tensor_generator.hpp:485
T operator()(Is...)
Definition: host_tensor_generator.hpp:489
std::normal_distribution< float > distribution
Definition: host_tensor_generator.hpp:483
Definition: host_tensor_generator.hpp:561
float operator()(Ts... Xs) const
Definition: host_tensor_generator.hpp:563
Definition: host_tensor_generator.hpp:654
T operator()(Ts... Xs) const
Definition: host_tensor_generator.hpp:658
T value
Definition: host_tensor_generator.hpp:655
ck::bf6x32_pk_t operator()(Ts... Xs) const
Definition: host_tensor_generator.hpp:639
ck::f4x2_pk_t operator()(Ts... Xs) const
Definition: host_tensor_generator.hpp:609
ck::f6x32_pk_t operator()(Ts... Xs) const
Definition: host_tensor_generator.hpp:622
Is used to generate sequential values based on the specified dimension.
Definition: host_tensor_generator.hpp:594
T operator()(Ts... Xs) const
Definition: host_tensor_generator.hpp:596
Unsigned representation of a conventional biased Float32 exponent.
Definition: e8m0.hpp:25
Definition: data_type.hpp:41
Definition: data_type.hpp:78
__host__ __device__ void pack(const T x, const index_t i)
Definition: data_type.hpp:118
Definition: data_type.hpp:186
Definition: functional2.hpp:33