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// Eviction priorities applicable for prefetch and applypriority intrinsics.
107enum class EvictPolicyType : uint8_t {
108 EVICT_NORMAL = 0, // default
109 EVICT_LAST = 1,
110};
111
112inline StringRef getEvictPolicyName(EvictPolicyType Policy) {
113 switch (Policy) {
114 case EvictPolicyType::EVICT_NORMAL:
115 return "L2::evict_normal";
116 case EvictPolicyType::EVICT_LAST:
117 return "L2::evict_last";
118 }
119 llvm_unreachable("invalid evict policy");
120}
121
122enum class Tcgen05MMAKind : uint8_t {
123 F16 = 0,
124 TF32 = 1,
125 F8F6F4 = 2,
126 I8 = 3,
127 TI16 = 4,
128};
129
136
137enum class Tcgen05MMACollectorBBuffer : uint8_t {
138 B0 = 0,
139 B1 = 1,
140 B2 = 2,
141 B3 = 3,
142};
143
144enum class TensormapElemType : uint8_t {
145 U8 = 0,
146 U16 = 1,
147 U32 = 2,
148 S32 = 3,
149 U64 = 4,
150 S64 = 5,
151 F16 = 6,
152 F32 = 7,
153 F32_FTZ = 8,
154 F64 = 9,
155 BF16 = 10,
156 TF32 = 11,
157 TF32_FTZ = 12,
158 B4x16 = 13,
159 B4x16_p64 = 14,
160 B6x16_p32 = 15,
161};
162
168
176
183
188
189// In-memory layout of an mbarrier object, as selected by the layout operand of
190// the llvm.nvvm.mbarrier.init and llvm.nvvm.mbarrier.check_layout intrinsics.
191enum class MBarrierLayout : uint8_t {
192 V0 = 0,
193 V1 = 1,
194};
195
196LLVM_ABI void printTcgen05MMAKind(raw_ostream &OS, const Constant *ImmArgVal);
197
198LLVM_ABI void printMBarrierLayout(raw_ostream &OS, const Constant *ImmArgVal);
199
200LLVM_ABI void printEvictPolicyType(raw_ostream &OS, const Constant *ImmArgVal);
201
202LLVM_ABI void printTMAReductionOp(raw_ostream &OS, const Constant *ImmArgVal);
203
204LLVM_ABI void printTMAValidateDataPattern(raw_ostream &OS,
205 const Constant *ImmArgVal);
206
207LLVM_ABI void printTcgen05CollectorUsageOp(raw_ostream &OS,
208 const Constant *ImmArgVal);
209
210LLVM_ABI void printTcgen05MMACollectorBBuffer(raw_ostream &OS,
211 const Constant *ImmArgVal);
212
213LLVM_ABI void printTensormapElemType(raw_ostream &OS,
214 const Constant *ImmArgVal);
215LLVM_ABI void printTensormapInterleaveLayout(raw_ostream &OS,
216 const Constant *ImmArgVal);
217LLVM_ABI void printTensormapSwizzleMode(raw_ostream &OS,
218 const Constant *ImmArgVal);
219LLVM_ABI void printTensormapSwizzleAtomicity(raw_ostream &OS,
220 const Constant *ImmArgVal);
221LLVM_ABI void printTensormapFillMode(raw_ostream &OS,
222 const Constant *ImmArgVal);
223LLVM_ABI void printFPRoundingMode(raw_ostream &OS, const Constant *ImmArgVal);
224
225inline bool FPToIntegerIntrinsicShouldFTZ(Intrinsic::ID IntrinsicID) {
226 switch (IntrinsicID) {
227 case Intrinsic::nvvm_f2i_rm_ftz:
228 case Intrinsic::nvvm_f2i_rn_ftz:
229 case Intrinsic::nvvm_f2i_rp_ftz:
230 case Intrinsic::nvvm_f2i_rz_ftz:
231
232 case Intrinsic::nvvm_f2ui_rm_ftz:
233 case Intrinsic::nvvm_f2ui_rn_ftz:
234 case Intrinsic::nvvm_f2ui_rp_ftz:
235 case Intrinsic::nvvm_f2ui_rz_ftz:
236
237 case Intrinsic::nvvm_f2ll_rm_ftz:
238 case Intrinsic::nvvm_f2ll_rn_ftz:
239 case Intrinsic::nvvm_f2ll_rp_ftz:
240 case Intrinsic::nvvm_f2ll_rz_ftz:
241
242 case Intrinsic::nvvm_f2ull_rm_ftz:
243 case Intrinsic::nvvm_f2ull_rn_ftz:
244 case Intrinsic::nvvm_f2ull_rp_ftz:
245 case Intrinsic::nvvm_f2ull_rz_ftz:
246 return true;
247
248 case Intrinsic::nvvm_f2i_rm:
249 case Intrinsic::nvvm_f2i_rn:
250 case Intrinsic::nvvm_f2i_rp:
251 case Intrinsic::nvvm_f2i_rz:
252
253 case Intrinsic::nvvm_f2ui_rm:
254 case Intrinsic::nvvm_f2ui_rn:
255 case Intrinsic::nvvm_f2ui_rp:
256 case Intrinsic::nvvm_f2ui_rz:
257
258 case Intrinsic::nvvm_d2i_rm:
259 case Intrinsic::nvvm_d2i_rn:
260 case Intrinsic::nvvm_d2i_rp:
261 case Intrinsic::nvvm_d2i_rz:
262
263 case Intrinsic::nvvm_d2ui_rm:
264 case Intrinsic::nvvm_d2ui_rn:
265 case Intrinsic::nvvm_d2ui_rp:
266 case Intrinsic::nvvm_d2ui_rz:
267
268 case Intrinsic::nvvm_f2ll_rm:
269 case Intrinsic::nvvm_f2ll_rn:
270 case Intrinsic::nvvm_f2ll_rp:
271 case Intrinsic::nvvm_f2ll_rz:
272
273 case Intrinsic::nvvm_f2ull_rm:
274 case Intrinsic::nvvm_f2ull_rn:
275 case Intrinsic::nvvm_f2ull_rp:
276 case Intrinsic::nvvm_f2ull_rz:
277
278 case Intrinsic::nvvm_d2ll_rm:
279 case Intrinsic::nvvm_d2ll_rn:
280 case Intrinsic::nvvm_d2ll_rp:
281 case Intrinsic::nvvm_d2ll_rz:
282
283 case Intrinsic::nvvm_d2ull_rm:
284 case Intrinsic::nvvm_d2ull_rn:
285 case Intrinsic::nvvm_d2ull_rp:
286 case Intrinsic::nvvm_d2ull_rz:
287 return false;
288 }
289 llvm_unreachable("Checking FTZ flag for invalid f2i/d2i intrinsic");
290}
291
292inline bool FPToIntegerIntrinsicResultIsSigned(Intrinsic::ID IntrinsicID) {
293 switch (IntrinsicID) {
294 // f2i
295 case Intrinsic::nvvm_f2i_rm:
296 case Intrinsic::nvvm_f2i_rm_ftz:
297 case Intrinsic::nvvm_f2i_rn:
298 case Intrinsic::nvvm_f2i_rn_ftz:
299 case Intrinsic::nvvm_f2i_rp:
300 case Intrinsic::nvvm_f2i_rp_ftz:
301 case Intrinsic::nvvm_f2i_rz:
302 case Intrinsic::nvvm_f2i_rz_ftz:
303 // d2i
304 case Intrinsic::nvvm_d2i_rm:
305 case Intrinsic::nvvm_d2i_rn:
306 case Intrinsic::nvvm_d2i_rp:
307 case Intrinsic::nvvm_d2i_rz:
308 // f2ll
309 case Intrinsic::nvvm_f2ll_rm:
310 case Intrinsic::nvvm_f2ll_rm_ftz:
311 case Intrinsic::nvvm_f2ll_rn:
312 case Intrinsic::nvvm_f2ll_rn_ftz:
313 case Intrinsic::nvvm_f2ll_rp:
314 case Intrinsic::nvvm_f2ll_rp_ftz:
315 case Intrinsic::nvvm_f2ll_rz:
316 case Intrinsic::nvvm_f2ll_rz_ftz:
317 // d2ll
318 case Intrinsic::nvvm_d2ll_rm:
319 case Intrinsic::nvvm_d2ll_rn:
320 case Intrinsic::nvvm_d2ll_rp:
321 case Intrinsic::nvvm_d2ll_rz:
322 return true;
323
324 // f2ui
325 case Intrinsic::nvvm_f2ui_rm:
326 case Intrinsic::nvvm_f2ui_rm_ftz:
327 case Intrinsic::nvvm_f2ui_rn:
328 case Intrinsic::nvvm_f2ui_rn_ftz:
329 case Intrinsic::nvvm_f2ui_rp:
330 case Intrinsic::nvvm_f2ui_rp_ftz:
331 case Intrinsic::nvvm_f2ui_rz:
332 case Intrinsic::nvvm_f2ui_rz_ftz:
333 // d2ui
334 case Intrinsic::nvvm_d2ui_rm:
335 case Intrinsic::nvvm_d2ui_rn:
336 case Intrinsic::nvvm_d2ui_rp:
337 case Intrinsic::nvvm_d2ui_rz:
338 // f2ull
339 case Intrinsic::nvvm_f2ull_rm:
340 case Intrinsic::nvvm_f2ull_rm_ftz:
341 case Intrinsic::nvvm_f2ull_rn:
342 case Intrinsic::nvvm_f2ull_rn_ftz:
343 case Intrinsic::nvvm_f2ull_rp:
344 case Intrinsic::nvvm_f2ull_rp_ftz:
345 case Intrinsic::nvvm_f2ull_rz:
346 case Intrinsic::nvvm_f2ull_rz_ftz:
347 // d2ull
348 case Intrinsic::nvvm_d2ull_rm:
349 case Intrinsic::nvvm_d2ull_rn:
350 case Intrinsic::nvvm_d2ull_rp:
351 case Intrinsic::nvvm_d2ull_rz:
352 return false;
353 }
354 llvm_unreachable(
355 "Checking invalid f2i/d2i intrinsic for signed int conversion");
356}
357
358inline bool FPToIntegerIntrinsicNaNZero(Intrinsic::ID IntrinsicID) {
359 switch (IntrinsicID) {
360 // f2i
361 case Intrinsic::nvvm_f2i_rm:
362 case Intrinsic::nvvm_f2i_rn:
363 case Intrinsic::nvvm_f2i_rp:
364 case Intrinsic::nvvm_f2i_rz:
365 case Intrinsic::nvvm_f2i_rm_ftz:
366 case Intrinsic::nvvm_f2i_rn_ftz:
367 case Intrinsic::nvvm_f2i_rp_ftz:
368 case Intrinsic::nvvm_f2i_rz_ftz:
369 // f2ui
370 case Intrinsic::nvvm_f2ui_rm:
371 case Intrinsic::nvvm_f2ui_rn:
372 case Intrinsic::nvvm_f2ui_rp:
373 case Intrinsic::nvvm_f2ui_rz:
374 case Intrinsic::nvvm_f2ui_rm_ftz:
375 case Intrinsic::nvvm_f2ui_rn_ftz:
376 case Intrinsic::nvvm_f2ui_rp_ftz:
377 case Intrinsic::nvvm_f2ui_rz_ftz:
378 return true;
379 // d2i
380 case Intrinsic::nvvm_d2i_rm:
381 case Intrinsic::nvvm_d2i_rn:
382 case Intrinsic::nvvm_d2i_rp:
383 case Intrinsic::nvvm_d2i_rz:
384 // d2ui
385 case Intrinsic::nvvm_d2ui_rm:
386 case Intrinsic::nvvm_d2ui_rn:
387 case Intrinsic::nvvm_d2ui_rp:
388 case Intrinsic::nvvm_d2ui_rz:
389 // f2ll
390 case Intrinsic::nvvm_f2ll_rm:
391 case Intrinsic::nvvm_f2ll_rn:
392 case Intrinsic::nvvm_f2ll_rp:
393 case Intrinsic::nvvm_f2ll_rz:
394 case Intrinsic::nvvm_f2ll_rm_ftz:
395 case Intrinsic::nvvm_f2ll_rn_ftz:
396 case Intrinsic::nvvm_f2ll_rp_ftz:
397 case Intrinsic::nvvm_f2ll_rz_ftz:
398 // f2ull
399 case Intrinsic::nvvm_f2ull_rm:
400 case Intrinsic::nvvm_f2ull_rn:
401 case Intrinsic::nvvm_f2ull_rp:
402 case Intrinsic::nvvm_f2ull_rz:
403 case Intrinsic::nvvm_f2ull_rm_ftz:
404 case Intrinsic::nvvm_f2ull_rn_ftz:
405 case Intrinsic::nvvm_f2ull_rp_ftz:
406 case Intrinsic::nvvm_f2ull_rz_ftz:
407 // d2ll
408 case Intrinsic::nvvm_d2ll_rm:
409 case Intrinsic::nvvm_d2ll_rn:
410 case Intrinsic::nvvm_d2ll_rp:
411 case Intrinsic::nvvm_d2ll_rz:
412 // d2ull
413 case Intrinsic::nvvm_d2ull_rm:
414 case Intrinsic::nvvm_d2ull_rn:
415 case Intrinsic::nvvm_d2ull_rp:
416 case Intrinsic::nvvm_d2ull_rz:
417 return false;
418 }
419 llvm_unreachable("Checking NaN result for invalid f2i/d2i intrinsic");
420}
421
422inline APFloat::roundingMode
423GetFPToIntegerRoundingMode(Intrinsic::ID IntrinsicID) {
424 switch (IntrinsicID) {
425 // RM:
426 case Intrinsic::nvvm_f2i_rm:
427 case Intrinsic::nvvm_f2ui_rm:
428 case Intrinsic::nvvm_f2i_rm_ftz:
429 case Intrinsic::nvvm_f2ui_rm_ftz:
430 case Intrinsic::nvvm_d2i_rm:
431 case Intrinsic::nvvm_d2ui_rm:
432
433 case Intrinsic::nvvm_f2ll_rm:
434 case Intrinsic::nvvm_f2ull_rm:
435 case Intrinsic::nvvm_f2ll_rm_ftz:
436 case Intrinsic::nvvm_f2ull_rm_ftz:
437 case Intrinsic::nvvm_d2ll_rm:
438 case Intrinsic::nvvm_d2ull_rm:
439 return APFloat::rmTowardNegative;
440
441 // RN:
442 case Intrinsic::nvvm_f2i_rn:
443 case Intrinsic::nvvm_f2ui_rn:
444 case Intrinsic::nvvm_f2i_rn_ftz:
445 case Intrinsic::nvvm_f2ui_rn_ftz:
446 case Intrinsic::nvvm_d2i_rn:
447 case Intrinsic::nvvm_d2ui_rn:
448
449 case Intrinsic::nvvm_f2ll_rn:
450 case Intrinsic::nvvm_f2ull_rn:
451 case Intrinsic::nvvm_f2ll_rn_ftz:
452 case Intrinsic::nvvm_f2ull_rn_ftz:
453 case Intrinsic::nvvm_d2ll_rn:
454 case Intrinsic::nvvm_d2ull_rn:
455 return APFloat::rmNearestTiesToEven;
456
457 // RP:
458 case Intrinsic::nvvm_f2i_rp:
459 case Intrinsic::nvvm_f2ui_rp:
460 case Intrinsic::nvvm_f2i_rp_ftz:
461 case Intrinsic::nvvm_f2ui_rp_ftz:
462 case Intrinsic::nvvm_d2i_rp:
463 case Intrinsic::nvvm_d2ui_rp:
464
465 case Intrinsic::nvvm_f2ll_rp:
466 case Intrinsic::nvvm_f2ull_rp:
467 case Intrinsic::nvvm_f2ll_rp_ftz:
468 case Intrinsic::nvvm_f2ull_rp_ftz:
469 case Intrinsic::nvvm_d2ll_rp:
470 case Intrinsic::nvvm_d2ull_rp:
471 return APFloat::rmTowardPositive;
472
473 // RZ:
474 case Intrinsic::nvvm_f2i_rz:
475 case Intrinsic::nvvm_f2ui_rz:
476 case Intrinsic::nvvm_f2i_rz_ftz:
477 case Intrinsic::nvvm_f2ui_rz_ftz:
478 case Intrinsic::nvvm_d2i_rz:
479 case Intrinsic::nvvm_d2ui_rz:
480
481 case Intrinsic::nvvm_f2ll_rz:
482 case Intrinsic::nvvm_f2ull_rz:
483 case Intrinsic::nvvm_f2ll_rz_ftz:
484 case Intrinsic::nvvm_f2ull_rz_ftz:
485 case Intrinsic::nvvm_d2ll_rz:
486 case Intrinsic::nvvm_d2ull_rz:
487 return APFloat::rmTowardZero;
488 }
489 llvm_unreachable("Checking rounding mode for invalid f2i/d2i intrinsic");
490}
491
492inline bool FMinFMaxShouldFTZ(Intrinsic::ID IntrinsicID) {
493 switch (IntrinsicID) {
494 case Intrinsic::nvvm_fmax_ftz_f:
495 case Intrinsic::nvvm_fmax_ftz_nan_f:
496 case Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_f:
497 case Intrinsic::nvvm_fmax_ftz_xorsign_abs_f:
498
499 case Intrinsic::nvvm_fmin_ftz_f:
500 case Intrinsic::nvvm_fmin_ftz_nan_f:
501 case Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_f:
502 case Intrinsic::nvvm_fmin_ftz_xorsign_abs_f:
503 return true;
504
505 case Intrinsic::nvvm_fmax_d:
506 case Intrinsic::nvvm_fmax_f:
507 case Intrinsic::nvvm_fmax_nan_f:
508 case Intrinsic::nvvm_fmax_nan_xorsign_abs_f:
509 case Intrinsic::nvvm_fmax_xorsign_abs_f:
510
511 case Intrinsic::nvvm_fmin_d:
512 case Intrinsic::nvvm_fmin_f:
513 case Intrinsic::nvvm_fmin_nan_f:
514 case Intrinsic::nvvm_fmin_nan_xorsign_abs_f:
515 case Intrinsic::nvvm_fmin_xorsign_abs_f:
516 return false;
517 }
518 llvm_unreachable("Checking FTZ flag for invalid fmin/fmax intrinsic");
519}
520
521inline bool FMinFMaxPropagatesNaNs(Intrinsic::ID IntrinsicID) {
522 switch (IntrinsicID) {
523 case Intrinsic::nvvm_fmax_ftz_nan_f:
524 case Intrinsic::nvvm_fmax_nan_f:
525 case Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_f:
526 case Intrinsic::nvvm_fmax_nan_xorsign_abs_f:
527
528 case Intrinsic::nvvm_fmin_ftz_nan_f:
529 case Intrinsic::nvvm_fmin_nan_f:
530 case Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_f:
531 case Intrinsic::nvvm_fmin_nan_xorsign_abs_f:
532 return true;
533
534 case Intrinsic::nvvm_fmax_d:
535 case Intrinsic::nvvm_fmax_f:
536 case Intrinsic::nvvm_fmax_ftz_f:
537 case Intrinsic::nvvm_fmax_ftz_xorsign_abs_f:
538 case Intrinsic::nvvm_fmax_xorsign_abs_f:
539
540 case Intrinsic::nvvm_fmin_d:
541 case Intrinsic::nvvm_fmin_f:
542 case Intrinsic::nvvm_fmin_ftz_f:
543 case Intrinsic::nvvm_fmin_ftz_xorsign_abs_f:
544 case Intrinsic::nvvm_fmin_xorsign_abs_f:
545 return false;
546 }
547 llvm_unreachable("Checking NaN flag for invalid fmin/fmax intrinsic");
548}
549
550inline bool FMinFMaxIsXorSignAbs(Intrinsic::ID IntrinsicID) {
551 switch (IntrinsicID) {
552 case Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_f:
553 case Intrinsic::nvvm_fmax_ftz_xorsign_abs_f:
554 case Intrinsic::nvvm_fmax_nan_xorsign_abs_f:
555 case Intrinsic::nvvm_fmax_xorsign_abs_f:
556
557 case Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_f:
558 case Intrinsic::nvvm_fmin_ftz_xorsign_abs_f:
559 case Intrinsic::nvvm_fmin_nan_xorsign_abs_f:
560 case Intrinsic::nvvm_fmin_xorsign_abs_f:
561 return true;
562
563 case Intrinsic::nvvm_fmax_d:
564 case Intrinsic::nvvm_fmax_f:
565 case Intrinsic::nvvm_fmax_ftz_f:
566 case Intrinsic::nvvm_fmax_ftz_nan_f:
567 case Intrinsic::nvvm_fmax_nan_f:
568
569 case Intrinsic::nvvm_fmin_d:
570 case Intrinsic::nvvm_fmin_f:
571 case Intrinsic::nvvm_fmin_ftz_f:
572 case Intrinsic::nvvm_fmin_ftz_nan_f:
573 case Intrinsic::nvvm_fmin_nan_f:
574 return false;
575 }
576 llvm_unreachable("Checking XorSignAbs flag for invalid fmin/fmax intrinsic");
577}
578
579inline bool UnaryMathIntrinsicShouldFTZ(Intrinsic::ID IntrinsicID) {
580 switch (IntrinsicID) {
581 case Intrinsic::nvvm_ceil_ftz_f:
582 case Intrinsic::nvvm_fabs_ftz:
583 case Intrinsic::nvvm_floor_ftz_f:
584 case Intrinsic::nvvm_round_ftz_f:
585 case Intrinsic::nvvm_saturate_ftz_f:
586 case Intrinsic::nvvm_sqrt_rn_ftz_f:
587 return true;
588 case Intrinsic::nvvm_ceil_f:
589 case Intrinsic::nvvm_ceil_d:
590 case Intrinsic::nvvm_fabs:
591 case Intrinsic::nvvm_floor_f:
592 case Intrinsic::nvvm_floor_d:
593 case Intrinsic::nvvm_round_f:
594 case Intrinsic::nvvm_round_d:
595 case Intrinsic::nvvm_saturate_d:
596 case Intrinsic::nvvm_saturate_f:
597 case Intrinsic::nvvm_sqrt_f:
598 case Intrinsic::nvvm_sqrt_rn_d:
599 case Intrinsic::nvvm_sqrt_rn_f:
600 return false;
601 }
602 llvm_unreachable("Checking FTZ flag for invalid unary intrinsic");
603}
604
605inline bool RCPShouldFTZ(Intrinsic::ID IntrinsicID) {
606 switch (IntrinsicID) {
607 case Intrinsic::nvvm_rcp_rm_ftz_f:
608 case Intrinsic::nvvm_rcp_rn_ftz_f:
609 case Intrinsic::nvvm_rcp_rp_ftz_f:
610 case Intrinsic::nvvm_rcp_rz_ftz_f:
611 return true;
612 case Intrinsic::nvvm_rcp_rm_d:
613 case Intrinsic::nvvm_rcp_rm_f:
614 case Intrinsic::nvvm_rcp_rn_d:
615 case Intrinsic::nvvm_rcp_rn_f:
616 case Intrinsic::nvvm_rcp_rp_d:
617 case Intrinsic::nvvm_rcp_rp_f:
618 case Intrinsic::nvvm_rcp_rz_d:
619 case Intrinsic::nvvm_rcp_rz_f:
620 return false;
621 }
622 llvm_unreachable("Checking FTZ flag for invalid rcp intrinsic");
623}
624
625inline APFloat::roundingMode GetRCPRoundingMode(Intrinsic::ID IntrinsicID) {
626 switch (IntrinsicID) {
627 case Intrinsic::nvvm_rcp_rm_f:
628 case Intrinsic::nvvm_rcp_rm_d:
629 case Intrinsic::nvvm_rcp_rm_ftz_f:
630 return APFloat::rmTowardNegative;
631
632 case Intrinsic::nvvm_rcp_rn_f:
633 case Intrinsic::nvvm_rcp_rn_d:
634 case Intrinsic::nvvm_rcp_rn_ftz_f:
635 return APFloat::rmNearestTiesToEven;
636
637 case Intrinsic::nvvm_rcp_rp_f:
638 case Intrinsic::nvvm_rcp_rp_d:
639 case Intrinsic::nvvm_rcp_rp_ftz_f:
640 return APFloat::rmTowardPositive;
641
642 case Intrinsic::nvvm_rcp_rz_f:
643 case Intrinsic::nvvm_rcp_rz_d:
644 case Intrinsic::nvvm_rcp_rz_ftz_f:
645 return APFloat::rmTowardZero;
646 }
647 llvm_unreachable("Checking rounding mode for invalid rcp intrinsic");
648}
649
650inline DenormalMode GetNVVMDenormMode(bool ShouldFTZ) {
651 if (ShouldFTZ)
652 return DenormalMode::getPreserveSign();
653 return DenormalMode::getIEEE();
654}
655
656inline APFloat::roundingMode GetRoundingModeFromImmArg(const Value *ImmArgVal) {
657 return static_cast<APFloat::roundingMode>(
658 cast<ConstantInt>(ImmArgVal)->getSExtValue());
659}
660
661inline StringRef GetRoundingModeName(APFloat::roundingMode RM) {
662 switch (RM) {
663 case APFloat::rmNearestTiesToEven:
664 return "rn";
665 case APFloat::rmTowardZero:
666 return "rz";
667 case APFloat::rmTowardNegative:
668 return "rm";
669 case APFloat::rmTowardPositive:
670 return "rp";
671 default:
672 return "";
673 }
674}
675
676inline bool FMulShouldFTZ(Intrinsic::ID IntrinsicID) {
677 switch (IntrinsicID) {
678 case Intrinsic::nvvm_mul_rm_ftz_f:
679 case Intrinsic::nvvm_mul_rn_ftz_f:
680 case Intrinsic::nvvm_mul_rp_ftz_f:
681 case Intrinsic::nvvm_mul_rz_ftz_f:
682 return true;
683
684 case Intrinsic::nvvm_mul_rm_f:
685 case Intrinsic::nvvm_mul_rn_f:
686 case Intrinsic::nvvm_mul_rp_f:
687 case Intrinsic::nvvm_mul_rz_f:
688 case Intrinsic::nvvm_mul_rm_d:
689 case Intrinsic::nvvm_mul_rn_d:
690 case Intrinsic::nvvm_mul_rp_d:
691 case Intrinsic::nvvm_mul_rz_d:
692 return false;
693 }
694 llvm_unreachable("Checking FTZ flag for invalid NVVM mul intrinsic");
695}
696
697inline APFloat::roundingMode GetFMulRoundingMode(Intrinsic::ID IntrinsicID) {
698 switch (IntrinsicID) {
699 case Intrinsic::nvvm_mul_rm_f:
700 case Intrinsic::nvvm_mul_rm_d:
701 case Intrinsic::nvvm_mul_rm_ftz_f:
702 return APFloat::rmTowardNegative;
703 case Intrinsic::nvvm_mul_rn_f:
704 case Intrinsic::nvvm_mul_rn_d:
705 case Intrinsic::nvvm_mul_rn_ftz_f:
706 return APFloat::rmNearestTiesToEven;
707 case Intrinsic::nvvm_mul_rp_f:
708 case Intrinsic::nvvm_mul_rp_d:
709 case Intrinsic::nvvm_mul_rp_ftz_f:
710 return APFloat::rmTowardPositive;
711 case Intrinsic::nvvm_mul_rz_f:
712 case Intrinsic::nvvm_mul_rz_d:
713 case Intrinsic::nvvm_mul_rz_ftz_f:
714 return APFloat::rmTowardZero;
715 }
716 llvm_unreachable("Invalid FP instrinsic rounding mode for NVVM mul");
717}
718
719inline bool FDivShouldFTZ(Intrinsic::ID IntrinsicID) {
720 switch (IntrinsicID) {
721 case Intrinsic::nvvm_div_rm_ftz_f:
722 case Intrinsic::nvvm_div_rn_ftz_f:
723 case Intrinsic::nvvm_div_rp_ftz_f:
724 case Intrinsic::nvvm_div_rz_ftz_f:
725 return true;
726
727 case Intrinsic::nvvm_div_rm_f:
728 case Intrinsic::nvvm_div_rn_f:
729 case Intrinsic::nvvm_div_rp_f:
730 case Intrinsic::nvvm_div_rz_f:
731 case Intrinsic::nvvm_div_rm_d:
732 case Intrinsic::nvvm_div_rn_d:
733 case Intrinsic::nvvm_div_rp_d:
734 case Intrinsic::nvvm_div_rz_d:
735 return false;
736 }
737 llvm_unreachable("Checking FTZ flag for invalid NVVM div intrinsic");
738}
739
740inline APFloat::roundingMode GetFDivRoundingMode(Intrinsic::ID IntrinsicID) {
741 switch (IntrinsicID) {
742 case Intrinsic::nvvm_div_rm_f:
743 case Intrinsic::nvvm_div_rm_d:
744 case Intrinsic::nvvm_div_rm_ftz_f:
745 return APFloat::rmTowardNegative;
746 case Intrinsic::nvvm_div_rn_f:
747 case Intrinsic::nvvm_div_rn_d:
748 case Intrinsic::nvvm_div_rn_ftz_f:
749 return APFloat::rmNearestTiesToEven;
750 case Intrinsic::nvvm_div_rp_f:
751 case Intrinsic::nvvm_div_rp_d:
752 case Intrinsic::nvvm_div_rp_ftz_f:
753 return APFloat::rmTowardPositive;
754 case Intrinsic::nvvm_div_rz_f:
755 case Intrinsic::nvvm_div_rz_d:
756 case Intrinsic::nvvm_div_rz_ftz_f:
757 return APFloat::rmTowardZero;
758 }
759 llvm_unreachable("Invalid FP instrinsic rounding mode for NVVM div");
760}
761
762inline bool FMAShouldFTZ(Intrinsic::ID IntrinsicID) {
763 switch (IntrinsicID) {
764 case Intrinsic::nvvm_fma_rm_ftz_f:
765 case Intrinsic::nvvm_fma_rn_ftz_f:
766 case Intrinsic::nvvm_fma_rp_ftz_f:
767 case Intrinsic::nvvm_fma_rz_ftz_f:
768 return true;
769
770 case Intrinsic::nvvm_fma_rm_f:
771 case Intrinsic::nvvm_fma_rn_f:
772 case Intrinsic::nvvm_fma_rp_f:
773 case Intrinsic::nvvm_fma_rz_f:
774 case Intrinsic::nvvm_fma_rm_d:
775 case Intrinsic::nvvm_fma_rn_d:
776 case Intrinsic::nvvm_fma_rp_d:
777 case Intrinsic::nvvm_fma_rz_d:
778 return false;
779 }
780 llvm_unreachable("Checking FTZ flag for invalid NVVM fma intrinsic");
781}
782
783inline APFloat::roundingMode GetFMARoundingMode(Intrinsic::ID IntrinsicID) {
784 switch (IntrinsicID) {
785 case Intrinsic::nvvm_fma_rm_f:
786 case Intrinsic::nvvm_fma_rm_d:
787 case Intrinsic::nvvm_fma_rm_ftz_f:
788 return APFloat::rmTowardNegative;
789 case Intrinsic::nvvm_fma_rn_f:
790 case Intrinsic::nvvm_fma_rn_d:
791 case Intrinsic::nvvm_fma_rn_ftz_f:
792 return APFloat::rmNearestTiesToEven;
793 case Intrinsic::nvvm_fma_rp_f:
794 case Intrinsic::nvvm_fma_rp_d:
795 case Intrinsic::nvvm_fma_rp_ftz_f:
796 return APFloat::rmTowardPositive;
797 case Intrinsic::nvvm_fma_rz_f:
798 case Intrinsic::nvvm_fma_rz_d:
799 case Intrinsic::nvvm_fma_rz_ftz_f:
800 return APFloat::rmTowardZero;
801 }
802 llvm_unreachable("Invalid FP instrinsic rounding mode for NVVM fma");
803}
804
805} // namespace nvvm
806} // namespace llvm
807#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:656
llvm::nvvm::GetFMARoundingMode
APFloat::roundingMode GetFMARoundingMode(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:783
llvm::nvvm::printTcgen05MMACollectorBBuffer
LLVM_ABI void printTcgen05MMACollectorBBuffer(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:109
llvm::nvvm::GetRoundingModeName
StringRef GetRoundingModeName(APFloat::roundingMode RM)
Definition NVVMIntrinsicUtils.h:661
llvm::nvvm::printTensormapSwizzleMode
LLVM_ABI void printTensormapSwizzleMode(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:162
llvm::nvvm::GetNVVMDenormMode
DenormalMode GetNVVMDenormMode(bool ShouldFTZ)
Definition NVVMIntrinsicUtils.h:650
llvm::nvvm::getEvictPolicyName
StringRef getEvictPolicyName(EvictPolicyType Policy)
Definition NVVMIntrinsicUtils.h:112
llvm::nvvm::printFPRoundingMode
LLVM_ABI void printFPRoundingMode(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:190
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:358
llvm::nvvm::printTensormapInterleaveLayout
LLVM_ABI void printTensormapInterleaveLayout(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:144
llvm::nvvm::GetFDivRoundingMode
APFloat::roundingMode GetFDivRoundingMode(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:740
llvm::nvvm::printTensormapSwizzleAtomicity
LLVM_ABI void printTensormapSwizzleAtomicity(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:176
llvm::nvvm::FPToIntegerIntrinsicResultIsSigned
bool FPToIntegerIntrinsicResultIsSigned(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:292
llvm::nvvm::GetFPToIntegerRoundingMode
APFloat::roundingMode GetFPToIntegerRoundingMode(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:423
llvm::nvvm::RCPShouldFTZ
bool RCPShouldFTZ(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:605
llvm::nvvm::getTMATensorReductionOpName
StringRef getTMATensorReductionOpName(TMAReductionOp Op)
Definition NVVMIntrinsicUtils.h:45
llvm::nvvm::FPToIntegerIntrinsicShouldFTZ
bool FPToIntegerIntrinsicShouldFTZ(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:225
llvm::nvvm::getTMAValidateDataPatternName
StringRef getTMAValidateDataPatternName(TMAValidateDataPattern Pattern)
Definition NVVMIntrinsicUtils.h:87
llvm::nvvm::FDivShouldFTZ
bool FDivShouldFTZ(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:719
llvm::nvvm::FMinFMaxIsXorSignAbs
bool FMinFMaxIsXorSignAbs(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:550
llvm::nvvm::GetFMulRoundingMode
APFloat::roundingMode GetFMulRoundingMode(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:697
llvm::nvvm::printTcgen05MMAKind
LLVM_ABI void printTcgen05MMAKind(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:65
llvm::nvvm::printTcgen05CollectorUsageOp
LLVM_ABI void printTcgen05CollectorUsageOp(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:88
llvm::nvvm::UnaryMathIntrinsicShouldFTZ
bool UnaryMathIntrinsicShouldFTZ(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:579
llvm::nvvm::FMinFMaxShouldFTZ
bool FMinFMaxShouldFTZ(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:492
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:762
llvm::nvvm::printMBarrierLayout
LLVM_ABI void printMBarrierLayout(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:52
llvm::nvvm::FMulShouldFTZ
bool FMulShouldFTZ(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:676
llvm::nvvm::printTensormapFillMode
LLVM_ABI void printTensormapFillMode(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:195
llvm::nvvm::GetRCPRoundingMode
APFloat::roundingMode GetRCPRoundingMode(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:625
llvm::nvvm::FMinFMaxPropagatesNaNs
bool FMinFMaxPropagatesNaNs(Intrinsic::ID IntrinsicID)
Definition NVVMIntrinsicUtils.h:521
llvm::nvvm::printTensormapElemType
LLVM_ABI void printTensormapElemType(raw_ostream &OS, const Constant *ImmArgVal)
Definition NVVMIntrinsicUtils.cpp:130
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