LLVM 24.0.0git
NVVMIntrinsicUtils.h
Go to the documentation of this file.
1//===--- NVVMIntrinsicUtils.h -----------------------------------*- C++ -*-===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8//
9/// \file
10/// This file contains the definitions of the enumerations and flags
11/// associated with NVVM Intrinsics, along with some helper functions.
12//
13//===----------------------------------------------------------------------===//
14
15#ifndef LLVM_IR_NVVMINTRINSICUTILS_H
16#define LLVM_IR_NVVMINTRINSICUTILS_H
17
18#include <stdint.h>
19
20#include "llvm/ADT/APFloat.h"
21#include "llvm/ADT/APInt.h"
22#include "llvm/ADT/StringRef.h"
23#include "llvm/IR/Constants.h"
24#include "llvm/IR/Intrinsics.h"
25#include "llvm/IR/IntrinsicsNVPTX.h"
26#include "llvm/Support/raw_ostream.h"
27
28namespace llvm {
29namespace nvvm {
30
31// Reduction Ops supported with TMA Copy from Shared
32// to Global Memory for the "cp.reduce.async.bulk.tensor.*"
33// family of PTX instructions.
34enum class TMAReductionOp : uint8_t {
35 ADD = 0,
36 MIN = 1,
37 MAX = 2,
38 INC = 3,
39 DEC = 4,
40 AND = 5,
41 OR = 6,
42 XOR = 7,
43};
44
45inline StringRef getTMATensorReductionOpName(TMAReductionOp Op) {
46 switch (Op) {
47 case TMAReductionOp::ADD:
48 return "add";
49 case TMAReductionOp::MIN:
50 return "min";
51 case TMAReductionOp::MAX:
52 return "max";
53 case TMAReductionOp::INC:
54 return "inc";
55 case TMAReductionOp::DEC:
56 return "dec";
57 case TMAReductionOp::AND:
58 return "and";
59 case TMAReductionOp::OR:
60 return "or";
61 case TMAReductionOp::XOR:
62 return "xor";
63 }
64 llvm_unreachable("invalid TMA tensorreduction operation");
65}
66
67// Enum to represent the cta_group::1 and
68// cta_group::2 variants in TMA/TCGEN05 family of
69// PTX instructions.
70enum class CTAGroupKind : uint8_t {
71 CG_NONE = 0, // default with no cta_group modifier
72 CG_1 = 1, // cta_group::1 modifier
73 CG_2 = 2, // cta_group::2 modifier
74};
75
76// Validate data patterns supported with the TMA Copy from Global to Shared
77// memory (CTA and Cluster) intrinsics.
78enum class TMAValidateDataPattern : uint8_t {
79 DISABLED = 0, // Default; no validate_data qualifier is emitted.
80 PER_16BYTES_80000000 = 1,
81 PER_16BYTES_8000 = 2,
82 PER_16BYTES_80 = 3,
83 PER_16BYTES_8 = 4,
84 PER_ELEMENT_FF = 5,
85};
86
87inline StringRef getTMAValidateDataPatternName(TMAValidateDataPattern Pattern) {
88 using VDTy = TMAValidateDataPattern;
89 switch (Pattern) {
90 case VDTy::DISABLED:
91 return "disabled";
92 case VDTy::PER_16BYTES_80000000:
93 return "per_16bytes::80000000";
94 case VDTy::PER_16BYTES_8000:
95 return "per_16bytes::8000";
96 case VDTy::PER_16BYTES_80:
97 return "per_16bytes::80";
98 case VDTy::PER_16BYTES_8:
99 return "per_16bytes::8";
100 case VDTy::PER_ELEMENT_FF:
101 return "per_element::ff";
102 }
103 llvm_unreachable("invalid TMA validate data pattern");
104}
105
106// Scope of the memory ordering semantics.
107enum class MemScope : uint8_t {
108 CTA = 0,
109 CLUSTER = 1,
110 GPU = 2,
111 SYS = 3,
112};
113
114inline StringRef getMemScopeName(MemScope Scope) {
115 switch (Scope) {
116 case MemScope::CTA:
117 return "cta";
118 case MemScope::CLUSTER:
119 return "cluster";
120 case MemScope::GPU:
121 return "gpu";
122 case MemScope::SYS:
123 return "sys";
124 }
125 llvm_unreachable("invalid memory scope");
126}
127
128// Eviction priorities applicable for prefetch and applypriority intrinsics.
129enum class EvictPolicyType : uint8_t {
130 EVICT_NORMAL = 0, // default
131 EVICT_LAST = 1,
132};
133
134inline StringRef getEvictPolicyName(EvictPolicyType Policy) {
135 switch (Policy) {
136 case EvictPolicyType::EVICT_NORMAL:
137 return "L2::evict_normal";
138 case EvictPolicyType::EVICT_LAST:
139 return "L2::evict_last";
140 }
141 llvm_unreachable("invalid evict policy");
142}
143
144enum class Tcgen05MMAKind : uint8_t {
145 F16 = 0,
146 TF32 = 1,
147 F8F6F4 = 2,
148 I8 = 3,
149 TI16 = 4,
150};
151
158
159enum class Tcgen05MMACollectorBBuffer : uint8_t {
160 B0 = 0,
161 B1 = 1,
162 B2 = 2,
163 B3 = 3,
164};
165
166enum class TensormapElemType : uint8_t {
167 U8 = 0,
168 U16 = 1,
169 U32 = 2,
170 S32 = 3,
171 U64 = 4,
172 S64 = 5,
173 F16 = 6,
174 F32 = 7,
175 F32_FTZ = 8,
176 F64 = 9,
177 BF16 = 10,
178 TF32 = 11,
179 TF32_FTZ = 12,
180 B4x16 = 13,
181 B4x16_p64 = 14,
182 B6x16_p32 = 15,
183};
184
190
198
205
210
211// In-memory layout of an mbarrier object, as selected by the layout operand of
212// the llvm.nvvm.mbarrier.init and llvm.nvvm.mbarrier.check_layout intrinsics.
213enum class MBarrierLayout : uint8_t {
214 V0 = 0,
215 V1 = 1,
216};
217
218LLVM_ABI void printTcgen05MMAKind(raw_ostream &OS, const Constant *ImmArgVal);
219
220LLVM_ABI void printMBarrierLayout(raw_ostream &OS, const Constant *ImmArgVal);
221
222LLVM_ABI void printEvictPolicyType(raw_ostream &OS, const Constant *ImmArgVal);
223
224LLVM_ABI void printTMAReductionOp(raw_ostream &OS, const Constant *ImmArgVal);
225
226LLVM_ABI void printTMAValidateDataPattern(raw_ostream &OS,
227 const Constant *ImmArgVal);
228
229LLVM_ABI void printMemScope(raw_ostream &OS, const Constant *ImmArgVal);
230
231LLVM_ABI void printTcgen05CollectorUsageOp(raw_ostream &OS,
232 const Constant *ImmArgVal);
233
234LLVM_ABI void printTcgen05MMACollectorBBuffer(raw_ostream &OS,
235 const Constant *ImmArgVal);
236
237LLVM_ABI void printTensormapElemType(raw_ostream &OS,
238 const Constant *ImmArgVal);
239LLVM_ABI void printTensormapInterleaveLayout(raw_ostream &OS,
240 const Constant *ImmArgVal);
241LLVM_ABI void printTensormapSwizzleMode(raw_ostream &OS,
242 const Constant *ImmArgVal);
243LLVM_ABI void printTensormapSwizzleAtomicity(raw_ostream &OS,
244 const Constant *ImmArgVal);
245LLVM_ABI void printTensormapFillMode(raw_ostream &OS,
246 const Constant *ImmArgVal);
247LLVM_ABI void printFPRoundingMode(raw_ostream &OS, const Constant *ImmArgVal);
248
249inline bool FPToIntegerIntrinsicShouldFTZ(Intrinsic::ID IntrinsicID) {
250 switch (IntrinsicID) {
251 case Intrinsic::nvvm_f2i_rm_ftz:
252 case Intrinsic::nvvm_f2i_rn_ftz:
253 case Intrinsic::nvvm_f2i_rp_ftz:
254 case Intrinsic::nvvm_f2i_rz_ftz:
255
256 case Intrinsic::nvvm_f2ui_rm_ftz:
257 case Intrinsic::nvvm_f2ui_rn_ftz:
258 case Intrinsic::nvvm_f2ui_rp_ftz:
259 case Intrinsic::nvvm_f2ui_rz_ftz:
260
261 case Intrinsic::nvvm_f2ll_rm_ftz:
262 case Intrinsic::nvvm_f2ll_rn_ftz:
263 case Intrinsic::nvvm_f2ll_rp_ftz:
264 case Intrinsic::nvvm_f2ll_rz_ftz:
265
266 case Intrinsic::nvvm_f2ull_rm_ftz:
267 case Intrinsic::nvvm_f2ull_rn_ftz:
268 case Intrinsic::nvvm_f2ull_rp_ftz:
269 case Intrinsic::nvvm_f2ull_rz_ftz:
270 return true;
271
272 case Intrinsic::nvvm_f2i_rm:
273 case Intrinsic::nvvm_f2i_rn:
274 case Intrinsic::nvvm_f2i_rp:
275 case Intrinsic::nvvm_f2i_rz:
276
277 case Intrinsic::nvvm_f2ui_rm:
278 case Intrinsic::nvvm_f2ui_rn:
279 case Intrinsic::nvvm_f2ui_rp:
280 case Intrinsic::nvvm_f2ui_rz:
281
282 case Intrinsic::nvvm_d2i_rm:
283 case Intrinsic::nvvm_d2i_rn:
284 case Intrinsic::nvvm_d2i_rp:
285 case Intrinsic::nvvm_d2i_rz:
286
287 case Intrinsic::nvvm_d2ui_rm:
288 case Intrinsic::nvvm_d2ui_rn:
289 case Intrinsic::nvvm_d2ui_rp:
290 case Intrinsic::nvvm_d2ui_rz:
291
292 case Intrinsic::nvvm_f2ll_rm:
293 case Intrinsic::nvvm_f2ll_rn:
294 case Intrinsic::nvvm_f2ll_rp:
295 case Intrinsic::nvvm_f2ll_rz:
296
297 case Intrinsic::nvvm_f2ull_rm:
298 case Intrinsic::nvvm_f2ull_rn:
299 case Intrinsic::nvvm_f2ull_rp:
300 case Intrinsic::nvvm_f2ull_rz:
301
302 case Intrinsic::nvvm_d2ll_rm:
303 case Intrinsic::nvvm_d2ll_rn:
304 case Intrinsic::nvvm_d2ll_rp:
305 case Intrinsic::nvvm_d2ll_rz:
306
307 case Intrinsic::nvvm_d2ull_rm:
308 case Intrinsic::nvvm_d2ull_rn:
309 case Intrinsic::nvvm_d2ull_rp:
310 case Intrinsic::nvvm_d2ull_rz:
311 return false;
312 }
313 llvm_unreachable("Checking FTZ flag for invalid f2i/d2i intrinsic");
314}
315
316inline bool FPToIntegerIntrinsicResultIsSigned(Intrinsic::ID IntrinsicID) {
317 switch (IntrinsicID) {
318 // f2i
319 case Intrinsic::nvvm_f2i_rm:
320 case Intrinsic::nvvm_f2i_rm_ftz:
321 case Intrinsic::nvvm_f2i_rn:
322 case Intrinsic::nvvm_f2i_rn_ftz:
323 case Intrinsic::nvvm_f2i_rp:
324 case Intrinsic::nvvm_f2i_rp_ftz:
325 case Intrinsic::nvvm_f2i_rz:
326 case Intrinsic::nvvm_f2i_rz_ftz:
327 // d2i
328 case Intrinsic::nvvm_d2i_rm:
329 case Intrinsic::nvvm_d2i_rn:
330 case Intrinsic::nvvm_d2i_rp:
331 case Intrinsic::nvvm_d2i_rz:
332 // f2ll
333 case Intrinsic::nvvm_f2ll_rm:
334 case Intrinsic::nvvm_f2ll_rm_ftz:
335 case Intrinsic::nvvm_f2ll_rn:
336 case Intrinsic::nvvm_f2ll_rn_ftz:
337 case Intrinsic::nvvm_f2ll_rp:
338 case Intrinsic::nvvm_f2ll_rp_ftz:
339 case Intrinsic::nvvm_f2ll_rz:
340 case Intrinsic::nvvm_f2ll_rz_ftz:
341 // d2ll
342 case Intrinsic::nvvm_d2ll_rm:
343 case Intrinsic::nvvm_d2ll_rn:
344 case Intrinsic::nvvm_d2ll_rp:
345 case Intrinsic::nvvm_d2ll_rz:
346 return true;
347
348 // f2ui
349 case Intrinsic::nvvm_f2ui_rm:
350 case Intrinsic::nvvm_f2ui_rm_ftz:
351 case Intrinsic::nvvm_f2ui_rn:
352 case Intrinsic::nvvm_f2ui_rn_ftz:
353 case Intrinsic::nvvm_f2ui_rp:
354 case Intrinsic::nvvm_f2ui_rp_ftz:
355 case Intrinsic::nvvm_f2ui_rz:
356 case Intrinsic::nvvm_f2ui_rz_ftz:
357 // d2ui
358 case Intrinsic::nvvm_d2ui_rm:
359 case Intrinsic::nvvm_d2ui_rn:
360 case Intrinsic::nvvm_d2ui_rp:
361 case Intrinsic::nvvm_d2ui_rz:
362 // f2ull
363 case Intrinsic::nvvm_f2ull_rm:
364 case Intrinsic::nvvm_f2ull_rm_ftz:
365 case Intrinsic::nvvm_f2ull_rn:
366 case Intrinsic::nvvm_f2ull_rn_ftz:
367 case Intrinsic::nvvm_f2ull_rp:
368 case Intrinsic::nvvm_f2ull_rp_ftz:
369 case Intrinsic::nvvm_f2ull_rz:
370 case Intrinsic::nvvm_f2ull_rz_ftz:
371 // d2ull
372 case Intrinsic::nvvm_d2ull_rm:
373 case Intrinsic::nvvm_d2ull_rn:
374 case Intrinsic::nvvm_d2ull_rp:
375 case Intrinsic::nvvm_d2ull_rz:
376 return false;
377 }
378 llvm_unreachable(
379 "Checking invalid f2i/d2i intrinsic for signed int conversion");
380}
381
382inline bool FPToIntegerIntrinsicNaNZero(Intrinsic::ID IntrinsicID) {
383 switch (IntrinsicID) {
384 // f2i
385 case Intrinsic::nvvm_f2i_rm:
386 case Intrinsic::nvvm_f2i_rn:
387 case Intrinsic::nvvm_f2i_rp:
388 case Intrinsic::nvvm_f2i_rz:
389 case Intrinsic::nvvm_f2i_rm_ftz:
390 case Intrinsic::nvvm_f2i_rn_ftz:
391 case Intrinsic::nvvm_f2i_rp_ftz:
392 case Intrinsic::nvvm_f2i_rz_ftz:
393 // f2ui
394 case Intrinsic::nvvm_f2ui_rm:
395 case Intrinsic::nvvm_f2ui_rn:
396 case Intrinsic::nvvm_f2ui_rp:
397 case Intrinsic::nvvm_f2ui_rz:
398 case Intrinsic::nvvm_f2ui_rm_ftz:
399 case Intrinsic::nvvm_f2ui_rn_ftz:
400 case Intrinsic::nvvm_f2ui_rp_ftz:
401 case Intrinsic::nvvm_f2ui_rz_ftz:
402 return true;
403 // d2i
404 case Intrinsic::nvvm_d2i_rm:
405 case Intrinsic::nvvm_d2i_rn:
406 case Intrinsic::nvvm_d2i_rp:
407 case Intrinsic::nvvm_d2i_rz:
408 // d2ui
409 case Intrinsic::nvvm_d2ui_rm:
410 case Intrinsic::nvvm_d2ui_rn:
411 case Intrinsic::nvvm_d2ui_rp:
412 case Intrinsic::nvvm_d2ui_rz:
413 // f2ll
414 case Intrinsic::nvvm_f2ll_rm:
415 case Intrinsic::nvvm_f2ll_rn:
416 case Intrinsic::nvvm_f2ll_rp:
417 case Intrinsic::nvvm_f2ll_rz:
418 case Intrinsic::nvvm_f2ll_rm_ftz:
419 case Intrinsic::nvvm_f2ll_rn_ftz:
420 case Intrinsic::nvvm_f2ll_rp_ftz:
421 case Intrinsic::nvvm_f2ll_rz_ftz:
422 // f2ull
423 case Intrinsic::nvvm_f2ull_rm:
424 case Intrinsic::nvvm_f2ull_rn:
425 case Intrinsic::nvvm_f2ull_rp:
426 case Intrinsic::nvvm_f2ull_rz:
427 case Intrinsic::nvvm_f2ull_rm_ftz:
428 case Intrinsic::nvvm_f2ull_rn_ftz:
429 case Intrinsic::nvvm_f2ull_rp_ftz:
430 case Intrinsic::nvvm_f2ull_rz_ftz:
431 // d2ll
432 case Intrinsic::nvvm_d2ll_rm:
433 case Intrinsic::nvvm_d2ll_rn:
434 case Intrinsic::nvvm_d2ll_rp:
435 case Intrinsic::nvvm_d2ll_rz:
436 // d2ull
437 case Intrinsic::nvvm_d2ull_rm:
438 case Intrinsic::nvvm_d2ull_rn:
439 case Intrinsic::nvvm_d2ull_rp:
440 case Intrinsic::nvvm_d2ull_rz:
441 return false;
442 }
443 llvm_unreachable("Checking NaN result for invalid f2i/d2i intrinsic");
444}
445
446inline APFloat::roundingMode
447GetFPToIntegerRoundingMode(Intrinsic::ID IntrinsicID) {
448 switch (IntrinsicID) {
449 // RM:
450 case Intrinsic::nvvm_f2i_rm:
451 case Intrinsic::nvvm_f2ui_rm:
452 case Intrinsic::nvvm_f2i_rm_ftz:
453 case Intrinsic::nvvm_f2ui_rm_ftz:
454 case Intrinsic::nvvm_d2i_rm:
455 case Intrinsic::nvvm_d2ui_rm:
456
457 case Intrinsic::nvvm_f2ll_rm:
458 case Intrinsic::nvvm_f2ull_rm:
459 case Intrinsic::nvvm_f2ll_rm_ftz:
460 case Intrinsic::nvvm_f2ull_rm_ftz:
461 case Intrinsic::nvvm_d2ll_rm:
462 case Intrinsic::nvvm_d2ull_rm:
463 return APFloat::rmTowardNegative;
464
465 // RN:
466 case Intrinsic::nvvm_f2i_rn:
467 case Intrinsic::nvvm_f2ui_rn:
468 case Intrinsic::nvvm_f2i_rn_ftz:
469 case Intrinsic::nvvm_f2ui_rn_ftz:
470 case Intrinsic::nvvm_d2i_rn:
471 case Intrinsic::nvvm_d2ui_rn:
472
473 case Intrinsic::nvvm_f2ll_rn:
474 case Intrinsic::nvvm_f2ull_rn:
475 case Intrinsic::nvvm_f2ll_rn_ftz:
476 case Intrinsic::nvvm_f2ull_rn_ftz:
477 case Intrinsic::nvvm_d2ll_rn:
478 case Intrinsic::nvvm_d2ull_rn:
479 return APFloat::rmNearestTiesToEven;
480
481 // RP:
482 case Intrinsic::nvvm_f2i_rp:
483 case Intrinsic::nvvm_f2ui_rp:
484 case Intrinsic::nvvm_f2i_rp_ftz:
485 case Intrinsic::nvvm_f2ui_rp_ftz:
486 case Intrinsic::nvvm_d2i_rp:
487 case Intrinsic::nvvm_d2ui_rp:
488
489 case Intrinsic::nvvm_f2ll_rp:
490 case Intrinsic::nvvm_f2ull_rp:
491 case Intrinsic::nvvm_f2ll_rp_ftz:
492 case Intrinsic::nvvm_f2ull_rp_ftz:
493 case Intrinsic::nvvm_d2ll_rp:
494 case Intrinsic::nvvm_d2ull_rp:
495 return APFloat::rmTowardPositive;
496
497 // RZ:
498 case Intrinsic::nvvm_f2i_rz:
499 case Intrinsic::nvvm_f2ui_rz:
500 case Intrinsic::nvvm_f2i_rz_ftz:
501 case Intrinsic::nvvm_f2ui_rz_ftz:
502 case Intrinsic::nvvm_d2i_rz:
503 case Intrinsic::nvvm_d2ui_rz:
504
505 case Intrinsic::nvvm_f2ll_rz:
506 case Intrinsic::nvvm_f2ull_rz:
507 case Intrinsic::nvvm_f2ll_rz_ftz:
508 case Intrinsic::nvvm_f2ull_rz_ftz:
509 case Intrinsic::nvvm_d2ll_rz:
510 case Intrinsic::nvvm_d2ull_rz:
511 return APFloat::rmTowardZero;
512 }
513 llvm_unreachable("Checking rounding mode for invalid f2i/d2i intrinsic");
514}
515
516inline bool FMinFMaxShouldFTZ(Intrinsic::ID IntrinsicID) {
517 switch (IntrinsicID) {
518 case Intrinsic::nvvm_fmax_ftz_f:
519 case Intrinsic::nvvm_fmax_ftz_nan_f:
520 case Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_f:
521 case Intrinsic::nvvm_fmax_ftz_xorsign_abs_f:
522
523 case Intrinsic::nvvm_fmin_ftz_f:
524 case Intrinsic::nvvm_fmin_ftz_nan_f:
525 case Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_f:
526 case Intrinsic::nvvm_fmin_ftz_xorsign_abs_f:
527 return true;
528
529 case Intrinsic::nvvm_fmax_d:
530 case Intrinsic::nvvm_fmax_f:
531 case Intrinsic::nvvm_fmax_nan_f:
532 case Intrinsic::nvvm_fmax_nan_xorsign_abs_f:
533 case Intrinsic::nvvm_fmax_xorsign_abs_f:
534
535 case Intrinsic::nvvm_fmin_d:
536 case Intrinsic::nvvm_fmin_f:
537 case Intrinsic::nvvm_fmin_nan_f:
538 case Intrinsic::nvvm_fmin_nan_xorsign_abs_f:
539 case Intrinsic::nvvm_fmin_xorsign_abs_f:
540 return false;
541 }
542 llvm_unreachable("Checking FTZ flag for invalid fmin/fmax intrinsic");
543}
544
545inline bool FMinFMaxPropagatesNaNs(Intrinsic::ID IntrinsicID) {
546 switch (IntrinsicID) {
547 case Intrinsic::nvvm_fmax_ftz_nan_f:
548 case Intrinsic::nvvm_fmax_nan_f:
549 case Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_f:
550 case Intrinsic::nvvm_fmax_nan_xorsign_abs_f:
551
552 case Intrinsic::nvvm_fmin_ftz_nan_f:
553 case Intrinsic::nvvm_fmin_nan_f:
554 case Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_f:
555 case Intrinsic::nvvm_fmin_nan_xorsign_abs_f:
556 return true;
557
558 case Intrinsic::nvvm_fmax_d:
559 case Intrinsic::nvvm_fmax_f:
560 case Intrinsic::nvvm_fmax_ftz_f:
561 case Intrinsic::nvvm_fmax_ftz_xorsign_abs_f:
562 case Intrinsic::nvvm_fmax_xorsign_abs_f:
563
564 case Intrinsic::nvvm_fmin_d:
565 case Intrinsic::nvvm_fmin_f:
566 case Intrinsic::nvvm_fmin_ftz_f:
567 case Intrinsic::nvvm_fmin_ftz_xorsign_abs_f:
568 case Intrinsic::nvvm_fmin_xorsign_abs_f:
569 return false;
570 }
571 llvm_unreachable("Checking NaN flag for invalid fmin/fmax intrinsic");
572}
573
574inline bool FMinFMaxIsXorSignAbs(Intrinsic::ID IntrinsicID) {
575 switch (IntrinsicID) {
576 case Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_f:
577 case Intrinsic::nvvm_fmax_ftz_xorsign_abs_f:
578 case Intrinsic::nvvm_fmax_nan_xorsign_abs_f:
579 case Intrinsic::nvvm_fmax_xorsign_abs_f:
580
581 case Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_f:
582 case Intrinsic::nvvm_fmin_ftz_xorsign_abs_f:
583 case Intrinsic::nvvm_fmin_nan_xorsign_abs_f:
584 case Intrinsic::nvvm_fmin_xorsign_abs_f:
585 return true;
586
587 case Intrinsic::nvvm_fmax_d:
588 case Intrinsic::nvvm_fmax_f:
589 case Intrinsic::nvvm_fmax_ftz_f:
590 case Intrinsic::nvvm_fmax_ftz_nan_f:
591 case Intrinsic::nvvm_fmax_nan_f:
592
593 case Intrinsic::nvvm_fmin_d:
594 case Intrinsic::nvvm_fmin_f:
595 case Intrinsic::nvvm_fmin_ftz_f:
596 case Intrinsic::nvvm_fmin_ftz_nan_f:
597 case Intrinsic::nvvm_fmin_nan_f:
598 return false;
599 }
600 llvm_unreachable("Checking XorSignAbs flag for invalid fmin/fmax intrinsic");
601}
602
603inline bool UnaryMathIntrinsicShouldFTZ(Intrinsic::ID IntrinsicID) {
604 switch (IntrinsicID) {
605 case Intrinsic::nvvm_ceil_ftz_f:
606 case Intrinsic::nvvm_fabs_ftz:
607 case Intrinsic::nvvm_floor_ftz_f:
608 case Intrinsic::nvvm_round_ftz_f:
609 case Intrinsic::nvvm_saturate_ftz_f:
610 case Intrinsic::nvvm_sqrt_rn_ftz_f:
611 return true;
612 case Intrinsic::nvvm_ceil_f:
613 case Intrinsic::nvvm_ceil_d:
614 case Intrinsic::nvvm_fabs:
615 case Intrinsic::nvvm_floor_f:
616 case Intrinsic::nvvm_floor_d:
617 case Intrinsic::nvvm_round_f:
618 case Intrinsic::nvvm_round_d:
619 case Intrinsic::nvvm_saturate_d:
620 case Intrinsic::nvvm_saturate_f:
621 case Intrinsic::nvvm_sqrt_f:
622 case Intrinsic::nvvm_sqrt_rn_d:
623 case Intrinsic::nvvm_sqrt_rn_f:
624 return false;
625 }
626 llvm_unreachable("Checking FTZ flag for invalid unary intrinsic");
627}
628
629inline bool RCPShouldFTZ(Intrinsic::ID IntrinsicID) {
630 switch (IntrinsicID) {
631 case Intrinsic::nvvm_rcp_rm_ftz_f:
632 case Intrinsic::nvvm_rcp_rn_ftz_f:
633 case Intrinsic::nvvm_rcp_rp_ftz_f:
634 case Intrinsic::nvvm_rcp_rz_ftz_f:
635 return true;
636 case Intrinsic::nvvm_rcp_rm_d:
637 case Intrinsic::nvvm_rcp_rm_f:
638 case Intrinsic::nvvm_rcp_rn_d:
639 case Intrinsic::nvvm_rcp_rn_f:
640 case Intrinsic::nvvm_rcp_rp_d:
641 case Intrinsic::nvvm_rcp_rp_f:
642 case Intrinsic::nvvm_rcp_rz_d:
643 case Intrinsic::nvvm_rcp_rz_f:
644 return false;
645 }
646 llvm_unreachable("Checking FTZ flag for invalid rcp intrinsic");
647}
648
649inline APFloat::roundingMode GetRCPRoundingMode(Intrinsic::ID IntrinsicID) {
650 switch (IntrinsicID) {
651 case Intrinsic::nvvm_rcp_rm_f:
652 case Intrinsic::nvvm_rcp_rm_d:
653 case Intrinsic::nvvm_rcp_rm_ftz_f:
654 return APFloat::rmTowardNegative;
655
656 case Intrinsic::nvvm_rcp_rn_f:
657 case Intrinsic::nvvm_rcp_rn_d:
658 case Intrinsic::nvvm_rcp_rn_ftz_f:
659 return APFloat::rmNearestTiesToEven;
660
661 case Intrinsic::nvvm_rcp_rp_f:
662 case Intrinsic::nvvm_rcp_rp_d:
663 case Intrinsic::nvvm_rcp_rp_ftz_f:
664 return APFloat::rmTowardPositive;
665
666 case Intrinsic::nvvm_rcp_rz_f:
667 case Intrinsic::nvvm_rcp_rz_d:
668 case Intrinsic::nvvm_rcp_rz_ftz_f:
669 return APFloat::rmTowardZero;
670 }
671 llvm_unreachable("Checking rounding mode for invalid rcp intrinsic");
672}
673
674inline DenormalMode GetNVVMDenormMode(bool ShouldFTZ) {
675 if (ShouldFTZ)
676 return DenormalMode::getPreserveSign();
677 return DenormalMode::getIEEE();
678}
679
680inline APFloat::roundingMode GetRoundingModeFromImmArg(const Value *ImmArgVal) {
681 return static_cast<APFloat::roundingMode>(
682 cast<ConstantInt>(ImmArgVal)->getSExtValue());
683}
684
685inline StringRef GetRoundingModeName(APFloat::roundingMode RM) {
686 switch (RM) {
687 case APFloat::rmNearestTiesToEven:
688 return "rn";
689 case APFloat::rmTowardZero:
690 return "rz";
691 case APFloat::rmTowardNegative:
692 return "rm";
693 case APFloat::rmTowardPositive:
694 return "rp";
695 default:
696 return "";
697 }
698}
699
700inline bool FMulShouldFTZ(Intrinsic::ID IntrinsicID) {
701 switch (IntrinsicID) {
702 case Intrinsic::nvvm_mul_rm_ftz_f:
703 case Intrinsic::nvvm_mul_rn_ftz_f:
704 case Intrinsic::nvvm_mul_rp_ftz_f:
705 case Intrinsic::nvvm_mul_rz_ftz_f:
706 return true;
707
708 case Intrinsic::nvvm_mul_rm_f:
709 case Intrinsic::nvvm_mul_rn_f:
710 case Intrinsic::nvvm_mul_rp_f:
711 case Intrinsic::nvvm_mul_rz_f:
712 case Intrinsic::nvvm_mul_rm_d:
713 case Intrinsic::nvvm_mul_rn_d:
714 case Intrinsic::nvvm_mul_rp_d:
715 case Intrinsic::nvvm_mul_rz_d:
716 return false;
717 }
718 llvm_unreachable("Checking FTZ flag for invalid NVVM mul intrinsic");
719}
720
721inline APFloat::roundingMode GetFMulRoundingMode(Intrinsic::ID IntrinsicID) {
722 switch (IntrinsicID) {
723 case Intrinsic::nvvm_mul_rm_f:
724 case Intrinsic::nvvm_mul_rm_d:
725 case Intrinsic::nvvm_mul_rm_ftz_f:
726 return APFloat::rmTowardNegative;
727 case Intrinsic::nvvm_mul_rn_f:
728 case Intrinsic::nvvm_mul_rn_d:
729 case Intrinsic::nvvm_mul_rn_ftz_f:
730 return APFloat::rmNearestTiesToEven;
731 case Intrinsic::nvvm_mul_rp_f:
732 case Intrinsic::nvvm_mul_rp_d:
733 case Intrinsic::nvvm_mul_rp_ftz_f:
734 return APFloat::rmTowardPositive;
735 case Intrinsic::nvvm_mul_rz_f:
736 case Intrinsic::nvvm_mul_rz_d:
737 case Intrinsic::nvvm_mul_rz_ftz_f:
738 return APFloat::rmTowardZero;
739 }
740 llvm_unreachable("Invalid FP instrinsic rounding mode for NVVM mul");
741}
742
743inline bool FDivShouldFTZ(Intrinsic::ID IntrinsicID) {
744 switch (IntrinsicID) {
745 case Intrinsic::nvvm_div_rm_ftz_f:
746 case Intrinsic::nvvm_div_rn_ftz_f:
747 case Intrinsic::nvvm_div_rp_ftz_f:
748 case Intrinsic::nvvm_div_rz_ftz_f:
749 return true;
750
751 case Intrinsic::nvvm_div_rm_f:
752 case Intrinsic::nvvm_div_rn_f:
753 case Intrinsic::nvvm_div_rp_f:
754 case Intrinsic::nvvm_div_rz_f:
755 case Intrinsic::nvvm_div_rm_d:
756 case Intrinsic::nvvm_div_rn_d:
757 case Intrinsic::nvvm_div_rp_d:
758 case Intrinsic::nvvm_div_rz_d:
759 return false;
760 }
761 llvm_unreachable("Checking FTZ flag for invalid NVVM div intrinsic");
762}
763
764inline APFloat::roundingMode GetFDivRoundingMode(Intrinsic::ID IntrinsicID) {
765 switch (IntrinsicID) {
766 case Intrinsic::nvvm_div_rm_f:
767 case Intrinsic::nvvm_div_rm_d:
768 case Intrinsic::nvvm_div_rm_ftz_f:
769 return APFloat::rmTowardNegative;
770 case Intrinsic::nvvm_div_rn_f:
771 case Intrinsic::nvvm_div_rn_d:
772 case Intrinsic::nvvm_div_rn_ftz_f:
773 return APFloat::rmNearestTiesToEven;
774 case Intrinsic::nvvm_div_rp_f:
775 case Intrinsic::nvvm_div_rp_d:
776 case Intrinsic::nvvm_div_rp_ftz_f:
777 return APFloat::rmTowardPositive;
778 case Intrinsic::nvvm_div_rz_f:
779 case Intrinsic::nvvm_div_rz_d:
780 case Intrinsic::nvvm_div_rz_ftz_f:
781 return APFloat::rmTowardZero;
782 }
783 llvm_unreachable("Invalid FP instrinsic rounding mode for NVVM div");
784}
785
786inline bool FMAShouldFTZ(Intrinsic::ID IntrinsicID) {
787 switch (IntrinsicID) {
788 case Intrinsic::nvvm_fma_rm_ftz_f:
789 case Intrinsic::nvvm_fma_rn_ftz_f:
790 case Intrinsic::nvvm_fma_rp_ftz_f:
791 case Intrinsic::nvvm_fma_rz_ftz_f:
792 return true;
793
794 case Intrinsic::nvvm_fma_rm_f:
795 case Intrinsic::nvvm_fma_rn_f:
796 case Intrinsic::nvvm_fma_rp_f:
797 case Intrinsic::nvvm_fma_rz_f:
798 case Intrinsic::nvvm_fma_rm_d:
799 case Intrinsic::nvvm_fma_rn_d:
800 case Intrinsic::nvvm_fma_rp_d:
801 case Intrinsic::nvvm_fma_rz_d:
802 return false;
803 }
804 llvm_unreachable("Checking FTZ flag for invalid NVVM fma intrinsic");
805}
806
807inline APFloat::roundingMode GetFMARoundingMode(Intrinsic::ID IntrinsicID) {
808 switch (IntrinsicID) {
809 case Intrinsic::nvvm_fma_rm_f:
810 case Intrinsic::nvvm_fma_rm_d:
811 case Intrinsic::nvvm_fma_rm_ftz_f:
812 return APFloat::rmTowardNegative;
813 case Intrinsic::nvvm_fma_rn_f:
814 case Intrinsic::nvvm_fma_rn_d:
815 case Intrinsic::nvvm_fma_rn_ftz_f:
816 return APFloat::rmNearestTiesToEven;
817 case Intrinsic::nvvm_fma_rp_f:
818 case Intrinsic::nvvm_fma_rp_d:
819 case Intrinsic::nvvm_fma_rp_ftz_f:
820 return APFloat::rmTowardPositive;
821 case Intrinsic::nvvm_fma_rz_f:
822 case Intrinsic::nvvm_fma_rz_d:
823 case Intrinsic::nvvm_fma_rz_ftz_f:
824 return APFloat::rmTowardZero;
825 }
826 llvm_unreachable("Invalid FP instrinsic rounding mode for NVVM fma");
827}
828
829} // namespace nvvm
830} // namespace llvm
831#endif // LLVM_IR_NVVMINTRINSICUTILS_H
APFloat.h
This file declares a class to represent arbitrary precision floating point values and provide a varie...
APInt.h
This file implements a class to represent arbitrary precision integral constant values and operations...
LLVM_ABI
#define LLVM_ABI
Definition Compiler.h:215
Constants.h
This file contains the declarations for the subclasses of Constant, which represent the different fla...
llvm::APFloatBase::rmTowardZero
static constexpr roundingMode rmTowardZero
Definition APFloat.h:365
llvm::APFloatBase::roundingMode
llvm::RoundingMode roundingMode
IEEE-754R 4.3: Rounding-direction attributes.
Definition APFloat.h:359
llvm::APFloatBase::rmTowardNegative
static constexpr roundingMode rmTowardNegative
Definition APFloat.h:364
llvm::APFloatBase::rmNearestTiesToEven
static constexpr roundingMode rmNearestTiesToEven
Definition APFloat.h:361
llvm::APFloatBase::rmTowardPositive
static constexpr roundingMode rmTowardPositive
Definition APFloat.h:363
llvm::Constant
This is an important base class in LLVM.
Definition Constant.h:43
llvm::StringRef
Represent a constant reference to a string, i.e.
Definition StringRef.h:56
llvm::Value
LLVM Value Representation.
Definition Value.h:75
llvm::raw_ostream
This class implements an extremely fast bulk output stream that can only output to a stream.
Definition raw_ostream.h:53
uint8_t
llvm_unreachable
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
Definition ErrorHandling.h:164
llvm::nvvm::GetRoundingModeFromImmArg
APFloat::roundingMode GetRoundingModeFromImmArg(const Value *ImmArgVal)
Definition NVVMIntrinsicUtils.h:680
llvm::nvvm::GetFMARoundingMode
APFloat::roundingMode GetFMARoundingMode(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:807
llvm::nvvm::printTcgen05MMACollectorBBuffer
LLVM_ABI void printTcgen05MMACollectorBBuffer(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:119
llvm::nvvm::GetRoundingModeName
StringRef GetRoundingModeName(APFloat::roundingMode RM)
Definition NVVMIntrinsicUtils.h:685
llvm::nvvm::printTensormapSwizzleMode
LLVM_ABI void printTensormapSwizzleMode(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:172
llvm::nvvm::GetNVVMDenormMode
DenormalMode GetNVVMDenormMode(bool ShouldFTZ)
Definition NVVMIntrinsicUtils.h:674
llvm::nvvm::getEvictPolicyName
StringRef getEvictPolicyName(EvictPolicyType Policy)
Definition NVVMIntrinsicUtils.h:134
llvm::nvvm::printFPRoundingMode
LLVM_ABI void printFPRoundingMode(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:200
llvm::nvvm::printTMAReductionOp
LLVM_ABI void printTMAReductionOp(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:29
llvm::nvvm::printTMAValidateDataPattern
LLVM_ABI void printTMAValidateDataPattern(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:39
llvm::nvvm::FPToIntegerIntrinsicNaNZero
bool FPToIntegerIntrinsicNaNZero(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:382
llvm::nvvm::printTensormapInterleaveLayout
LLVM_ABI void printTensormapInterleaveLayout(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:154
llvm::nvvm::GetFDivRoundingMode
APFloat::roundingMode GetFDivRoundingMode(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:764
llvm::nvvm::printTensormapSwizzleAtomicity
LLVM_ABI void printTensormapSwizzleAtomicity(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:186
llvm::nvvm::FPToIntegerIntrinsicResultIsSigned
bool FPToIntegerIntrinsicResultIsSigned(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:316
llvm::nvvm::GetFPToIntegerRoundingMode
APFloat::roundingMode GetFPToIntegerRoundingMode(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:447
llvm::nvvm::RCPShouldFTZ
bool RCPShouldFTZ(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:629
llvm::nvvm::getTMATensorReductionOpName
StringRef getTMATensorReductionOpName(TMAReductionOp Op)
Definition NVVMIntrinsicUtils.h:45
llvm::nvvm::FPToIntegerIntrinsicShouldFTZ
bool FPToIntegerIntrinsicShouldFTZ(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:249
llvm::nvvm::printMemScope
LLVM_ABI void printMemScope(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:65
llvm::nvvm::getTMAValidateDataPatternName
StringRef getTMAValidateDataPatternName(TMAValidateDataPattern Pattern)
Definition NVVMIntrinsicUtils.h:87
llvm::nvvm::FDivShouldFTZ
bool FDivShouldFTZ(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:743
llvm::nvvm::FMinFMaxIsXorSignAbs
bool FMinFMaxIsXorSignAbs(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:574
llvm::nvvm::GetFMulRoundingMode
APFloat::roundingMode GetFMulRoundingMode(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:721
llvm::nvvm::printTcgen05MMAKind
LLVM_ABI void printTcgen05MMAKind(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:75
llvm::nvvm::printTcgen05CollectorUsageOp
LLVM_ABI void printTcgen05CollectorUsageOp(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:98
llvm::nvvm::UnaryMathIntrinsicShouldFTZ
bool UnaryMathIntrinsicShouldFTZ(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:603
llvm::nvvm::FMinFMaxShouldFTZ
bool FMinFMaxShouldFTZ(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:516
llvm::nvvm::printEvictPolicyType
LLVM_ABI void printEvictPolicyType(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:19
llvm::nvvm::FMAShouldFTZ
bool FMAShouldFTZ(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:786
llvm::nvvm::printMBarrierLayout
LLVM_ABI void printMBarrierLayout(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:52
llvm::nvvm::getMemScopeName
StringRef getMemScopeName(MemScope Scope)
Definition NVVMIntrinsicUtils.h:114
llvm::nvvm::FMulShouldFTZ
bool FMulShouldFTZ(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:700
llvm::nvvm::printTensormapFillMode
LLVM_ABI void printTensormapFillMode(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:205
llvm::nvvm::GetRCPRoundingMode
APFloat::roundingMode GetRCPRoundingMode(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:649
llvm::nvvm::FMinFMaxPropagatesNaNs
bool FMinFMaxPropagatesNaNs(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:545
llvm::nvvm::printTensormapElemType
LLVM_ABI void printTensormapElemType(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:140
llvm
This is an optimization pass for GlobalISel generic memory operations.
Definition FunctionInfo.h:25
llvm::Op
DWARFExpression::Operation Op
Definition DWARFExpressionPrinter.cpp:24
llvm::cast
decltype(auto) cast(const From &Val)
cast<X> - Return the argument parameter cast to the specified type.
Definition Casting.h:559
raw_ostream.h
llvm::DenormalMode
Represent subnormal handling kind for floating point instruction inputs and outputs.
Definition FloatingPointMode.h:71
llvm::DenormalMode::getPreserveSign
static constexpr DenormalMode getPreserveSign()
Definition FloatingPointMode.h:119
llvm::DenormalMode::getIEEE
static constexpr DenormalMode getIEEE()
Definition FloatingPointMode.h:115
Generated on for LLVM by doxygen 1.14.0