host_tensor_generator.hpp Source File

host_tensor_generator.hpp Source File#

Composable Kernel: 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
12template <typename T>
14{
15 template <typename... Is>
16 T operator()(Is...)
17 {
18 return T{0};
19 }
20};
21
22template <typename T>
24{
25 T value = 1;
26
27 template <typename... Is>
28 T operator()(Is...)
29 {
30 return value;
31 }
32};
33
34template <>
35struct GeneratorTensor_1<ck::half_t>
36{
37 float value = 1.0;
38
39 template <typename... Is>
44};
45
46template <>
47struct GeneratorTensor_1<ck::bhalf_t>
48{
49 float value = 1.0;
50
51 template <typename... Is>
56};
57
58#if defined CK_ENABLE_FP8
59template <>
61{
62 float value = 1.0;
63
64 template <typename... Is>
65 ck::f8_t operator()(Is...)
66 {
68 }
69};
70
71template <>
73{
74 float value = 1.0;
75
76 template <typename... Is>
78 {
80 }
81};
82#endif
83
84template <>
85struct GeneratorTensor_1<ck::f4_t>
86{
87 float value = 1.0;
88
89 template <typename... Is>
94};
95
96template <>
97struct GeneratorTensor_1<ck::f4x2_pk_t>
98{
99 float value = 1.0;
100
101 template <typename... Is>
106};
107
108template <>
109struct GeneratorTensor_1<ck::f6x32_pk_t>
110{
111 float value = 1.0;
112
113 template <typename... Is>
115 {
117 ck::static_for<0, 32, 1>{}([&](auto i) {
119 });
120 return r;
121 }
122};
123
124template <>
125struct GeneratorTensor_1<ck::bf6x32_pk_t>
126{
127 float value = 1.0;
128
129 template <typename... Is>
131 {
133 ck::static_for<0, 32, 1>{}([&](auto i) {
135 });
136 return r;
137 }
138};
139
140template <>
142{
144
145 template <typename... Is>
147 {
148 return value;
149 }
150};
151
152template <>
153struct 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
166template <>
167struct GeneratorTensor_1<ck::e8m0_bexp_t>
168{
169 float value = 1;
170
171 template <typename... Is>
176};
177
178template <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
191template <>
192struct GeneratorTensor_2<ck::f6x32_pk_t>
193{
194 int min_value = 0;
195 int max_value = 1;
196
197 template <typename... Is>
199 {
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
210template <>
211struct GeneratorTensor_2<ck::bf6x32_pk_t>
212{
213 int min_value = 0;
214 int max_value = 1;
215
216 template <typename... Is>
218 {
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
229template <>
230struct GeneratorTensor_2<ck::bhalf_t>
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;
240 }
241};
242
243template <>
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
256template <>
257struct 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
273template <>
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
289template <>
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
304template <>
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
318template <>
319struct 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;
330 }
331};
332
333template <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
348template <>
349struct GeneratorTensor_3<ck::bhalf_t>
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
366template <>
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
385template <>
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
403template <>
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
420template <>
421struct 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
439template <>
440struct 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
455template <>
456struct GeneratorTensor_3<ck::f6x32_pk_t>
457{
458 float min_value = 0;
459 float max_value = 1;
460
461 template <typename... Is>
463 {
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
475template <>
476struct GeneratorTensor_3<ck::bf6x32_pk_t>
477{
478 float min_value = 0;
479 float max_value = 1;
480
481 template <typename... Is>
483 {
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
495template <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
513template <>
514struct 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
532template <>
533struct GeneratorTensor_4<ck::f6x32_pk_t>
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 {
545 ck::static_for<0, 32, 1>{}([&](auto i) {
547 static_cast<ck::index_t>(i));
548 });
549
550 return r;
551 }
552};
553
554template <>
555struct GeneratorTensor_4<ck::bf6x32_pk_t>
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 {
567 ck::static_for<0, 32, 1>{}([&](auto i) {
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
608template <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
621template <ck::index_t Dim>
622struct 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];
631 }
632};
633
634template <ck::index_t Dim>
635struct GeneratorTensor_Sequential<ck::f6x32_pk_t, Dim>
636{
637 template <typename... Ts>
639 {
640 std::array<ck::index_t, sizeof...(Ts)> dims = {{static_cast<ck::index_t>(Xs)...}};
641
642 float tmp = dims[Dim];
643
646 [&](auto i) { r.pack(ck::type_convert<ck::f6_t>(tmp), static_cast<ck::index_t>(i)); });
647 return r;
648 }
649};
650
651template <ck::index_t Dim>
652struct GeneratorTensor_Sequential<ck::bf6x32_pk_t, Dim>
653{
654 template <typename... Ts>
656 {
657 std::array<ck::index_t, sizeof...(Ts)> dims = {{static_cast<ck::index_t>(Xs)...}};
658
659 float tmp = dims[Dim];
660
663 [&](auto i) { r.pack(ck::type_convert<ck::bf6_t>(tmp), static_cast<ck::index_t>(i)); });
664 return r;
665 }
666};
667
668template <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
ushort bhalf_t
Definition data_type.hpp:30
f8_fnuz_t f8_t
Definition amd_ck_fp8.hpp:1762
int32_t index_t
Definition ck.hpp:299
_Float16 half_t
Definition data_type.hpp:31
f6_pk_t< bf6_t, 32 > bf6x32_pk_t
Definition data_type.hpp:183
typename vector_type< float, 2 >::type float2_t
Definition dtype_vector.hpp:2145
unsigned _BitInt(4) f4_t
Definition data_type.hpp:33
bf8_fnuz_t bf8_t
Definition amd_ck_fp8.hpp:1763
__host__ __device__ constexpr Y type_convert(X x)
Definition utility/type_convert.hpp:98
f6_pk_t< f6_t, 32 > f6x32_pk_t
Definition data_type.hpp:181
signed char int8_t
Definition stdint.h:121
Definition host_tensor_generator.hpp:14
T operator()(Is...)
Definition host_tensor_generator.hpp:16
float value
Definition host_tensor_generator.hpp:127
ck::bf6x32_pk_t operator()(Is...)
Definition host_tensor_generator.hpp:130
ck::bhalf_t operator()(Is...)
Definition host_tensor_generator.hpp:52
float value
Definition host_tensor_generator.hpp:49
ck::e8m0_bexp_t operator()(Is...)
Definition host_tensor_generator.hpp:172
float value
Definition host_tensor_generator.hpp:169
float value
Definition host_tensor_generator.hpp:87
ck::f4_t operator()(Is...)
Definition host_tensor_generator.hpp:90
ck::f4x2_pk_t operator()(Is...)
Definition host_tensor_generator.hpp:102
float value
Definition host_tensor_generator.hpp:99
float value
Definition host_tensor_generator.hpp:111
ck::f6x32_pk_t operator()(Is...)
Definition host_tensor_generator.hpp:114
float value
Definition host_tensor_generator.hpp:37
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 value
Definition host_tensor_generator.hpp:155
int8_t operator()(Is...)
Definition host_tensor_generator.hpp:146
int8_t value
Definition host_tensor_generator.hpp:143
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
int min_value
Definition host_tensor_generator.hpp:213
int max_value
Definition host_tensor_generator.hpp:214
int max_value
Definition host_tensor_generator.hpp:233
int min_value
Definition host_tensor_generator.hpp:232
ck::bhalf_t operator()(Is...)
Definition host_tensor_generator.hpp:236
int min_value
Definition host_tensor_generator.hpp:307
ck::f4_t operator()(Is...)
Definition host_tensor_generator.hpp:311
int max_value
Definition host_tensor_generator.hpp:308
int max_value
Definition host_tensor_generator.hpp:322
int min_value
Definition host_tensor_generator.hpp:321
ck::f4x2_pk_t operator()(Is...)
Definition host_tensor_generator.hpp:325
int min_value
Definition host_tensor_generator.hpp:194
ck::f6x32_pk_t operator()(Is...)
Definition host_tensor_generator.hpp:198
int max_value
Definition host_tensor_generator.hpp:195
ck::pk_i4_t operator()(Is...)
Definition host_tensor_generator.hpp:263
int max_value
Definition host_tensor_generator.hpp:260
int min_value
Definition host_tensor_generator.hpp:259
int8_t operator()(Is...)
Definition host_tensor_generator.hpp:250
int min_value
Definition host_tensor_generator.hpp:246
int max_value
Definition host_tensor_generator.hpp:247
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
float min_value
Definition host_tensor_generator.hpp:478
float max_value
Definition host_tensor_generator.hpp:479
ck::bf6x32_pk_t operator()(Is...)
Definition host_tensor_generator.hpp:482
ck::bhalf_t operator()(Is...)
Definition host_tensor_generator.hpp:355
float min_value
Definition host_tensor_generator.hpp:351
float max_value
Definition host_tensor_generator.hpp:352
float min_value
Definition host_tensor_generator.hpp:406
ck::f4_t operator()(Is...)
Definition host_tensor_generator.hpp:410
float max_value
Definition host_tensor_generator.hpp:407
float max_value
Definition host_tensor_generator.hpp:424
float min_value
Definition host_tensor_generator.hpp:423
ck::f4x2_pk_t operator()(Is...)
Definition host_tensor_generator.hpp:427
ck::f6x32_pk_t operator()(Is...)
Definition host_tensor_generator.hpp:462
float min_value
Definition host_tensor_generator.hpp:458
float max_value
Definition host_tensor_generator.hpp:459
int max_value
Definition host_tensor_generator.hpp:443
int min_value
Definition host_tensor_generator.hpp:442
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
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
Unsigned representation of a conventional biased Float32 exponent.
Definition utility/e8m0.hpp:26
Definition data_type.hpp:42
__host__ __device__ void pack(const T x, const index_t i)
Definition data_type.hpp:119
Definition data_type.hpp:187
Definition functional2.hpp:33