[mlir][llvm] Ensure immediate usage in intrinsics

[Dinistro]

Summary

This commit changes intrinsics that have immarg parameter attributes to
model these parameters as attributes, instead of operands. Using
operands only works if the operation is an llvm.mlir.constant,
otherwise the exported LLVMIR is invalid.

[gysit]

I have only some minor comments.

@ftynse it may make sense you have a look. The change potentially impacts downstream projects.

mlir/test/Conversion/MemRefToLLVM/convert-dynamic-memref-ops.mlir

301

nit: remaining

mlir/test/Target/LLVMIR/llvmir-intrinsics.mlir

466

Should we have an export test with i32 as well?

[Dinistro]

address review comments

[gysit]

Thanks! LGTM

flang/lib/Optimizer/Builder/IntrinsicCall.cpp

 }
 
 mlir::Value IntrinsicLibrary::genIsFPClass(mlir::Type resultType,
                                            llvm::ArrayRef<mlir::Value> args,
                                            int fpclass) {
   assert(args.size() == 1);
   mlir::MLIRContext *context = builder.getContext();
   mlir::IntegerType i1ty = mlir::IntegerType::get(context, 1);
-  mlir::IntegerType i32ty = mlir::IntegerType::get(context, 32);
 
-  mlir::Value test = builder.createIntegerConstant(loc, i32ty, fpclass);
   mlir::Value isfpclass =
-      builder.create<mlir::LLVM::IsFPClass>(loc, i1ty, args[0], test);
+      builder.create<mlir::LLVM::IsFPClass>(loc, i1ty, args[0], fpclass);
   return builder.createConvert(loc, resultType, isfpclass);
 }
 
 mlir::Value IntrinsicLibrary::genIsNan(mlir::Type resultType,
                                        llvm::ArrayRef<mlir::Value> args) {
   // Check is signaling or quiet nan
   return genIsFPClass(resultType, args, 0b11);
 }

flang/test/Lower/Intrinsics/ieee_is_normal.f90

 ! RUN: bbc -emit-fir %s -o - | FileCheck %s
 ! RUN: flang-new -fc1 -emit-fir %s -o - | FileCheck %s
 
 ! CHECK-LABEL: ieee_is_normal_f16
 subroutine ieee_is_normal_f16(r)
   use ieee_arithmetic
   real(KIND=2) :: r
   i = ieee_is_normal(r)
-  ! CHECK: %[[test:.*]] = arith.constant 360 : i32
-  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}, %[[test]]) : (f16, i32) -> i1
+  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}) <{bit = 360 : i32}> : (f16) -> i1
   ! CHECK: fir.convert %[[l]] : (i1) -> !fir.logical<4>
 end subroutine ieee_is_normal_f16
 
 ! CHECK-LABEL: ieee_is_normal_bf16
 subroutine ieee_is_normal_bf16(r)
   use ieee_arithmetic
   real(KIND=3) :: r
   i = ieee_is_normal(r)
-  ! CHECK: %[[test:.*]] = arith.constant 360 : i32
-  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}, %[[test]]) : (bf16, i32) -> i1
+  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}) <{bit = 360 : i32}> : (bf16) -> i1
   ! CHECK: fir.convert %[[l]] : (i1) -> !fir.logical<4>
 end subroutine ieee_is_normal_bf16
 
 
 
 ! CHECK-LABEL: ieee_is_normal_f32
 subroutine ieee_is_normal_f32(r)
   use ieee_arithmetic
   real :: r
   i = ieee_is_normal(r)
-  ! CHECK: %[[test:.*]] = arith.constant 360 : i32
-  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}, %[[test]]) : (f32, i32) -> i1
+  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}) <{bit = 360 : i32}> : (f32) -> i1
   ! CHECK: fir.convert %[[l]] : (i1) -> !fir.logical<4>
 end subroutine ieee_is_normal_f32
 
 ! CHECK-LABEL: ieee_is_normal_f64
 subroutine ieee_is_normal_f64(r)
   use ieee_arithmetic
   real(KIND=8) :: r
   i = ieee_is_normal(r)
-  ! CHECK: %[[test:.*]] = arith.constant 360 : i32
-  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}, %[[test]]) : (f64, i32) -> i1
+  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}) <{bit = 360 : i32}> : (f64) -> i1
   ! CHECK: fir.convert %[[l]] : (i1) -> !fir.logical<4>
 end subroutine ieee_is_normal_f64
 
 ! CHECK-LABEL: ieee_is_normal_f80
 subroutine ieee_is_normal_f80(r)
   use ieee_arithmetic
   real(KIND=10) :: r
   i = ieee_is_normal(r)
-  ! CHECK: %[[test:.*]] = arith.constant 360 : i32
-  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}, %[[test]]) : (f80, i32) -> i1
+  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}) <{bit = 360 : i32}> : (f80) -> i1
   ! CHECK: fir.convert %[[l]] : (i1) -> !fir.logical<4>
 end subroutine ieee_is_normal_f80
 
 ! CHECK-LABEL: ieee_is_normal_f128
 subroutine ieee_is_normal_f128(r)
   use ieee_arithmetic
   real(KIND=16) :: r
   i = ieee_is_normal(r)
-  ! CHECK: %[[test:.*]] = arith.constant 360 : i32
-  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}, %[[test]]) : (f128, i32) -> i1
+  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}) <{bit = 360 : i32}> : (f128) -> i1
   ! CHECK: fir.convert %[[l]] : (i1) -> !fir.logical<4>
 end subroutine ieee_is_normal_f128

flang/test/Lower/Intrinsics/isnan.f90

 ! RUN: bbc -emit-fir %s -o - | FileCheck %s
 ! RUN: flang-new -fc1 -emit-fir %s -o - | FileCheck %s
 
 ! CHECK-LABEL: isnan_f32
 subroutine isnan_f32(r)
   real :: r
   i = isnan(r)
-  ! CHECK: %[[test:.*]] = arith.constant 3 : i32
-  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}, %[[test]]) : (f32, i32) -> i1
+  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}) <{bit = 3 : i32}> : (f32) -> i1
   ! CHECK: fir.convert %[[l]] : (i1) -> !fir.logical<4>
 end subroutine isnan_f32
 
 ! CHECK-LABEL: ieee_is_nan_f32
 subroutine ieee_is_nan_f32(r)
   use ieee_arithmetic
   real :: r
   i = ieee_is_nan(r)
-  ! CHECK: %[[test:.*]] = arith.constant 3 : i32
-  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}, %[[test]]) : (f32, i32) -> i1
+  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}) <{bit = 3 : i32}> : (f32) -> i1
   ! CHECK: fir.convert %[[l]] : (i1) -> !fir.logical<4>
 end subroutine ieee_is_nan_f32
 
 ! CHECK-LABEL: isnan_f64
 subroutine isnan_f64(r)
   real(KIND=8) :: r
   i = isnan(r)
-  ! CHECK: %[[test:.*]] = arith.constant 3 : i32
-  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}, %[[test]]) : (f64, i32) -> i1
+  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}) <{bit = 3 : i32}> : (f64) -> i1
   ! CHECK: fir.convert %[[l]] : (i1) -> !fir.logical<4>
 end subroutine isnan_f64
 
 ! CHECK-LABEL: ieee_is_nan_f64
 subroutine ieee_is_nan_f64(r)
   use ieee_arithmetic
   real(KIND=8) :: r
   i = ieee_is_nan(r)
-  ! CHECK: %[[test:.*]] = arith.constant 3 : i32
-  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}, %[[test]]) : (f64, i32) -> i1
+  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}) <{bit = 3 : i32}> : (f64) -> i1
   ! CHECK: fir.convert %[[l]] : (i1) -> !fir.logical<4>
 end subroutine ieee_is_nan_f64
 
 ! CHECK-LABEL: isnan_f80
 subroutine isnan_f80(r)
   real(KIND=10) :: r
   i = isnan(r)
-  ! CHECK: %[[test:.*]] = arith.constant 3 : i32
-  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}, %[[test]]) : (f80, i32) -> i1
+  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}) <{bit = 3 : i32}> : (f80) -> i1
   ! CHECK: fir.convert %[[l]] : (i1) -> !fir.logical<4>
 end subroutine isnan_f80
 
 ! CHECK-LABEL: ieee_is_nan_f80
 subroutine ieee_is_nan_f80(r)
   use ieee_arithmetic
   real(KIND=10) :: r
   i = ieee_is_nan(r)
-  ! CHECK: %[[test:.*]] = arith.constant 3 : i32
-  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}, %[[test]]) : (f80, i32) -> i1
+  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}) <{bit = 3 : i32}> : (f80) -> i1
   ! CHECK: fir.convert %[[l]] : (i1) -> !fir.logical<4>
 end subroutine ieee_is_nan_f80
 
 ! CHECK-LABEL: isnan_f128
 subroutine isnan_f128(r)
   real(KIND=16) :: r
   i = isnan(r)
-  ! CHECK: %[[test:.*]] = arith.constant 3 : i32
-  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}, %[[test]]) : (f128, i32) -> i1
+  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}) <{bit = 3 : i32}> : (f128) -> i1
   ! CHECK: fir.convert %[[l]] : (i1) -> !fir.logical<4>
 end subroutine isnan_f128
 
 ! CHECK-LABEL: ieee_is_nan_f128
 subroutine ieee_is_nan_f128(r)
   use ieee_arithmetic
   real(KIND=16) :: r
   i = ieee_is_nan(r)
-  ! CHECK: %[[test:.*]] = arith.constant 3 : i32
-  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}, %[[test]]) : (f128, i32) -> i1
+  ! CHECK: %[[l:.*]] = "llvm.intr.is.fpclass"(%{{.*}}) <{bit = 3 : i32}> : (f128) -> i1
   ! CHECK: fir.convert %[[l]] : (i1) -> !fir.logical<4>
 end subroutine ieee_is_nan_f128

mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td

dag fmfArg = (
     ins DefaultValuedAttr<LLVM_FastmathFlagsAttr, "{}">:$fastmathFlags);
   let arguments = !con(commonArgs, fmfArg);
 }
 
 class LLVM_CountZerosIntrOp<string func, list<Trait> traits = []> :
     LLVM_OneResultIntrOp<func, [], [0],
            !listconcat([Pure], traits)> {
   let arguments = (ins LLVM_ScalarOrVectorOf<AnySignlessInteger>:$in,
-                   I1:$zero_undefined);
+                   I1Attr:$is_zero_poison);
+  string mlirBuilder = [{
+    auto op = $_builder.create<$_qualCppClassName>($_location,
+      $_resultType, $in, $_int_attr($is_zero_poison));
+      $res = op;
+  }];
+  string llvmBuilder = [{
+    auto *inst = createIntrinsicCall(
+      builder, llvm::Intrinsic::}] # llvmEnumName # [{,
+      {$in, builder.getInt1(op.getIsZeroPoison())}, }]
+      # declTypes # [{);
+    $res = inst;
+  }];
 }
 
 def LLVM_AbsOp : LLVM_OneResultIntrOp<"abs", [], [0], [Pure]> {
   let arguments = (ins LLVM_ScalarOrVectorOf<AnySignlessInteger>:$in,
-                   I1:$is_int_min_poison);
+                   I1Attr:$is_int_min_poison);
+  string mlirBuilder = [{
+    auto op = $_builder.create<$_qualCppClassName>($_location,
+      $_resultType, $in, $_int_attr($is_int_min_poison));
+      $res = op;
+  }];
+  string llvmBuilder = [{
+    auto *inst = createIntrinsicCall(
+      builder, llvm::Intrinsic::}] # llvmEnumName # [{,
+      {$in, builder.getInt1(op.getIsIntMinPoison())}, }]
+      # declTypes # [{);
+    $res = inst;
+  }];
 }
 
 def LLVM_IsFPClass : LLVM_OneResultIntrOp<"is.fpclass", [], [0], [Pure]> {
-  let arguments = (ins LLVM_ScalarOrVectorOf<LLVM_AnyFloat>:$in, I32:$bit);
+  let arguments = (ins LLVM_ScalarOrVectorOf<LLVM_AnyFloat>:$in, I32Attr:$bit);
+  string mlirBuilder = [{
+    auto op = $_builder.create<$_qualCppClassName>($_location,
+      $_resultType, $in, $_int_attr($bit));
+      $res = op;
+  }];
+  string llvmBuilder = [{
+    auto *inst = createIntrinsicCall(
+      builder, llvm::Intrinsic::}] # llvmEnumName # [{,
+      {$in, builder.getInt32(op.getBit())},
+      }] # declTypes # [{);
+    $res = inst;
+  }];
 }
 
 def LLVM_CopySignOp : LLVM_BinarySameArgsIntrOpF<"copysign">;
 def LLVM_CosOp : LLVM_UnaryIntrOpF<"cos">;
 def LLVM_ExpOp : LLVM_UnaryIntrOpF<"exp">;
 def LLVM_Exp2Op : LLVM_UnaryIntrOpF<"exp2">;
 def LLVM_FAbsOp : LLVM_UnaryIntrOpF<"fabs">;
 def LLVM_FCeilOp : LLVM_UnaryIntrOpF<"ceil">;
 def LLVM_FFloorOp : LLVM_UnaryIntrOpF<"floor">;
 def LLVM_FMAOp : LLVM_TernarySameArgsIntrOpF<"fma">;
 def LLVM_FMulAddOp : LLVM_TernarySameArgsIntrOpF<"fmuladd">;
 def LLVM_Log10Op : LLVM_UnaryIntrOpF<"log10">;
 def LLVM_Log2Op : LLVM_UnaryIntrOpF<"log2">;
 def LLVM_LogOp : LLVM_UnaryIntrOpF<"log">;
 def LLVM_Prefetch : LLVM_ZeroResultIntrOp<"prefetch", [0]> {
-  let arguments = (ins LLVM_AnyPointer:$addr, I32:$rw, I32:$hint, I32:$cache);
+  let arguments = (ins LLVM_AnyPointer:$addr, I32Attr:$rw, I32Attr:$hint, I32Attr:$cache);
+  string mlirBuilder = [{
+    $_op = $_builder.create<$_qualCppClassName>($_location,
+      $addr, $_int_attr($rw), $_int_attr($hint), $_int_attr($cache));
+  }];
+  string llvmBuilder = [{
+    createIntrinsicCall(
+      builder, llvm::Intrinsic::}] # llvmEnumName # [{,
+      {$addr, builder.getInt32(op.getRw()),
+       builder.getInt32(op.getHint()),
+       builder.getInt32(op.getCache())},
+      }] # declTypes # [{);
+  }];
 }
 def LLVM_SinOp : LLVM_UnaryIntrOpF<"sin">;
 def LLVM_RoundEvenOp : LLVM_UnaryIntrOpF<"roundeven">;
 def LLVM_RoundOp : LLVM_UnaryIntrOpF<"round">;
 def LLVM_FTruncOp : LLVM_UnaryIntrOpF<"trunc">;
 def LLVM_SqrtOp : LLVM_UnaryIntrOpF<"sqrt">;
 def LLVM_PowOp : LLVM_BinarySameArgsIntrOpF<"pow">;
 def LLVM_PowIOp : LLVM_OneResultIntrOp<"powi", [], [0,1],

 def LLVM_UMaxOp : LLVM_BinarySameArgsIntrOpI<"umax">;
 def LLVM_UMinOp : LLVM_BinarySameArgsIntrOpI<"umin">;
 
 class LLVM_MemcpyIntrOpBase<string name> :
     LLVM_ZeroResultIntrOp<name, [0, 1, 2], [], /*requiresAccessGroup=*/1,
                                                /*requiresAliasAnalysis=*/1> {
   dag args = (ins Arg<LLVM_AnyPointer,"",[MemWrite]>:$dst,
                   Arg<LLVM_AnyPointer,"",[MemRead]>:$src,
-                  AnySignlessInteger:$len, I1:$isVolatile);
+                  AnySignlessInteger:$len, I1Attr:$isVolatile);
   // Append the alias attributes defined by LLVM_IntrOpBase.
   let arguments = !con(args, aliasAttrs);
   let builders = [
     OpBuilder<(ins "Value":$dst, "Value":$src, "Value":$len,
-                   "Value":$isVolatile), [{
+                   "bool":$isVolatile), [{
+      build($_builder, $_state, dst, src, len,
+            $_builder.getBoolAttr(isVolatile));
+    }]>,
+    OpBuilder<(ins "Value":$dst, "Value":$src, "Value":$len,
+                   "IntegerAttr":$isVolatile), [{
       build($_builder, $_state, dst, src, len, isVolatile,
             /*access_groups=*/nullptr, /*alias_scopes=*/nullptr,
             /*noalias_scopes=*/nullptr, /*tbaa=*/nullptr);
-    }]
-  >];
+    }]>
+  ];
+  string mlirBuilder = [{
+    $_op = $_builder.create<$_qualCppClassName>($_location,
+      $dst, $src, $len, $_int_attr($isVolatile));
+  }];
+  string llvmBuilder = [{
+    auto *inst = createIntrinsicCall(
+      builder, llvm::Intrinsic::}] # llvmEnumName # [{,
+      {$dst, $src, $len,
+       builder.getInt1(op.getIsVolatile())},
+      }] # declTypes # [{ ); }]
+      # setAccessGroupsMetadataCode
+      # setAliasAnalysisMetadataCode;
 }
 
 def LLVM_MemcpyOp : LLVM_MemcpyIntrOpBase<"memcpy">;
-def LLVM_MemcpyInlineOp : LLVM_MemcpyIntrOpBase<"memcpy.inline">;
 def LLVM_MemmoveOp : LLVM_MemcpyIntrOpBase<"memmove">;
 
+def LLVM_MemcpyInlineOp :
+    LLVM_ZeroResultIntrOp<"memcpy.inline", [0, 1], [],
+                          /*requiresAccessGroup=*/1,
+                          /*requiresAliasAnalysis=*/1> {
+  dag args = (ins Arg<LLVM_AnyPointer,"",[MemWrite]>:$dst,
+                  Arg<LLVM_AnyPointer,"",[MemRead]>:$src,
+                  APIntAttr:$len, I1Attr:$isVolatile);
+  // Append the alias attributes defined by LLVM_IntrOpBase.
+  let arguments = !con(args, aliasAttrs);
+  let builders = [
+    OpBuilder<(ins "Value":$dst, "Value":$src, "IntegerAttr":$len,
+                   "bool":$isVolatile), [{
+      build($_builder, $_state, dst, src, len,
+            $_builder.getBoolAttr(isVolatile));
+    }]>,
+    OpBuilder<(ins "Value":$dst, "Value":$src, "IntegerAttr":$len,
+                   "IntegerAttr":$isVolatile), [{
+      build($_builder, $_state, dst, src, len, isVolatile,
+            /*access_groups=*/nullptr, /*alias_scopes=*/nullptr,
+            /*noalias_scopes=*/nullptr, /*tbaa=*/nullptr);
+    }]>
+  ];
+  string mlirBuilder = [{
+    $_op = $_builder.create<$_qualCppClassName>($_location,
+      $dst, $src, $_int_attr($len), $_int_attr($isVolatile));
+  }];
+  string llvmBuilder = [{
+    auto *inst = createIntrinsicCall(
+      builder, llvm::Intrinsic::}] # llvmEnumName # [{,
+      {$dst, $src, builder.getInt(op.getLen()),
+       builder.getInt1(op.getIsVolatile())}, { }]
+      # !interleave(!listconcat(declTypeList, [
+       [{ moduleTranslation.convertType(op.getLenAttr().getType()) }]
+       ]), ", ") # [{ }); }]
+      # setAccessGroupsMetadataCode
+      # setAliasAnalysisMetadataCode;
+}
+
 def LLVM_MemsetOp : LLVM_ZeroResultIntrOp<"memset", [0, 2], [],
                       /*requiresAccessGroup=*/1, /*requiresAliasAnalysis=*/1> {
   dag args = (ins Arg<LLVM_AnyPointer,"",[MemWrite]>:$dst,
-                  I8:$val, AnySignlessInteger:$len, I1:$isVolatile);
+                  I8:$val, AnySignlessInteger:$len, I1Attr:$isVolatile);
   // Append the alias attributes defined by LLVM_IntrOpBase.
   let arguments = !con(args, aliasAttrs);
   let builders = [
     OpBuilder<(ins "Value":$dst, "Value":$val, "Value":$len,
-                    "Value":$isVolatile), [{
+                    "bool":$isVolatile), [{
+      build($_builder, $_state, dst, val, len,
+            $_builder.getBoolAttr(isVolatile));
+    }]>,
+    OpBuilder<(ins "Value":$dst, "Value":$val, "Value":$len,
+                    "IntegerAttr":$isVolatile), [{
       build($_builder, $_state, dst, val, len, isVolatile,
             /*access_groups=*/nullptr, /*alias_scopes=*/nullptr,
             /*noalias_scopes=*/nullptr, /*tbaa=*/nullptr);
-    }]
-  >];
+    }]>
+  ];
+  string mlirBuilder = [{
+    $_op = $_builder.create<$_qualCppClassName>($_location,
+      $dst, $val, $len, $_int_attr($isVolatile));
+  }];
+  string llvmBuilder = [{
+    auto *inst = createIntrinsicCall(
+      builder, llvm::Intrinsic::}] # llvmEnumName # [{,
+      {$dst, $val, $len,
+       builder.getInt1(op.getIsVolatile())},
+      }] # declTypes # [{ ); }]
+      # setAccessGroupsMetadataCode
+      # setAliasAnalysisMetadataCode;
 }
 
 def LLVM_NoAliasScopeDeclOp
     : LLVM_ZeroResultIntrOp<"experimental.noalias.scope.decl"> {
   let arguments = (ins SymbolRefAttr:$scope);
   string llvmBuilder = [{
     // Wrap the scope argument into a list since the LLVM IR intrinsic takes
     // a list containing exactly one scope rather than a scope itself.

mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td

dag aliasAttrs = !con(
             (ins OptionalAttr<SymbolRefArrayAttr>:$alias_scopes,
                  OptionalAttr<SymbolRefArrayAttr>:$noalias_scopes,
                  OptionalAttr<SymbolRefArrayAttr>:$tbaa),
             (ins )));
   string resultPattern = !if(!gt(numResults, 1),
                              LLVM_IntrPatterns.structResult,
                              LLVM_IntrPatterns.result);
   string llvmEnumName = enumName;
+  list<string> declTypeList = !listconcat(
+            ListIntSubst<resultPattern, overloadedResults>.lst,
+            ListIntSubst<LLVM_IntrPatterns.operand,
+                         overloadedOperands>.lst);
+  string declTypes = [{ { }] # !interleave(declTypeList, ", ") # [{ } }];
   let llvmBuilder = [{
     llvm::Module *module = builder.GetInsertBlock()->getModule();
     llvm::Function *fn = llvm::Intrinsic::getDeclaration(
         module,
-        llvm::Intrinsic::}] # enumName # [{,
-        { }] # !interleave(!listconcat(
-            ListIntSubst<resultPattern, overloadedResults>.lst,
-            ListIntSubst<LLVM_IntrPatterns.operand,
-                         overloadedOperands>.lst), ", ") # [{
-        });
+        llvm::Intrinsic::}] # enumName # [{,}] # declTypes # [{);
     auto operands = moduleTranslation.lookupValues(opInst.getOperands());
     }] # [{
     auto *inst = builder.CreateCall(fn, operands);
     (void) inst;
     }] # !if(!gt(requiresAccessGroup, 0), setAccessGroupsMetadataCode, "")
        # !if(!gt(requiresAliasAnalysis, 0), setAliasAnalysisMetadataCode, "")
        # !if(!gt(numResults, 0), "$res = inst;", "");
 

mlir/include/mlir/Target/LLVMIR/ModuleTranslation.h

 /// attributes and combinations thereof. Also, an array attribute with two
 /// elements is supported to represent a complex constant.  In case of error,
 /// report it to `loc` and return nullptr.
 llvm::Constant *getLLVMConstant(llvm::Type *llvmType, Attribute attr,
                                 Location loc,
                                 const ModuleTranslation &moduleTranslation);
 
 /// Creates a call to an LLVM IR intrinsic function with the given arguments.
-llvm::Value *createIntrinsicCall(llvm::IRBuilderBase &builder,
+llvm::CallInst *createIntrinsicCall(llvm::IRBuilderBase &builder,
                                     llvm::Intrinsic::ID intrinsic,
                                     ArrayRef<llvm::Value *> args = {},
                                     ArrayRef<llvm::Type *> tys = {});
 } // namespace detail
 
 } // namespace LLVM
 } // namespace mlir
 
 namespace llvm {
 template <typename T>
 struct isa_impl<T, ::mlir::LLVM::ModuleTranslation::StackFrame> {
   static inline bool
   doit(const ::mlir::LLVM::ModuleTranslation::StackFrame &frame) {
     return frame.getTypeID() == ::mlir::TypeID::get<T>();
   }
 };
 } // namespace llvm
 
 #endif // MLIR_TARGET_LLVMIR_MODULETRANSLATION_H

mlir/lib/Conversion/LLVMCommon/Pattern.cpp

LogicalResult ConvertToLLVMPattern::copyUnrankedDescriptors(
 
   // Compute allocation sizes.
   SmallVector<Value> sizes;
   UnrankedMemRefDescriptor::computeSizes(builder, loc, *getTypeConverter(),
                                          unrankedMemrefs, unrankedAddressSpaces,
                                          sizes);
 
   // Get frequently used types.
-  MLIRContext *context = builder.getContext();
-  auto i1Type = IntegerType::get(context, 1);
   Type indexType = getTypeConverter()->getIndexType();
 
   // Find the malloc and free, or declare them if necessary.
   auto module = builder.getInsertionPoint()->getParentOfType<ModuleOp>();
   LLVM::LLVMFuncOp freeFunc, mallocFunc;
   if (toDynamic)
     mallocFunc = LLVM::lookupOrCreateMallocFn(
         module, indexType, getTypeConverter()->useOpaquePointers());
   if (!toDynamic)
     freeFunc = LLVM::lookupOrCreateFreeFn(
         module, getTypeConverter()->useOpaquePointers());
 
-  // Initialize shared constants.
-  Value zero =
-      builder.create<LLVM::ConstantOp>(loc, i1Type, builder.getBoolAttr(false));
-
   unsigned unrankedMemrefPos = 0;
   for (unsigned i = 0, e = operands.size(); i < e; ++i) {
     Type type = origTypes[i];
     if (!isa<UnrankedMemRefType>(type))
       continue;
     Value allocationSize = sizes[unrankedMemrefPos++];
     UnrankedMemRefDescriptor desc(operands[i]);
 
     // Allocate memory, copy, and free the source if necessary.
     Value memory =
         toDynamic
             ? builder.create<LLVM::CallOp>(loc, mallocFunc, allocationSize)
                   .getResult()
             : builder.create<LLVM::AllocaOp>(loc, getVoidPtrType(),
                                              IntegerType::get(getContext(), 8),
                                              allocationSize,
                                              /*alignment=*/0);
     Value source = desc.memRefDescPtr(builder, loc);
-    builder.create<LLVM::MemcpyOp>(loc, memory, source, allocationSize, zero);
+    builder.create<LLVM::MemcpyOp>(loc, memory, source, allocationSize, false);
     if (!toDynamic)
       builder.create<LLVM::CallOp>(loc, freeFunc, source);
 
     // Create a new descriptor. The same descriptor can be returned multiple
     // times, attempting to modify its pointer can lead to memory leaks
     // (allocated twice and overwritten) or double frees (the caller does not
     // know if the descriptor points to the same memory).
     Type descriptorType = getTypeConverter()->convertType(type);

mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp

matchAndRewrite(MathOp op, typename MathOp::Adaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto operandType = adaptor.getOperand().getType();
 
     if (!operandType || !LLVM::isCompatibleType(operandType))
       return failure();
 
     auto loc = op.getLoc();
     auto resultType = op.getResult().getType();
-    auto boolZero = rewriter.getBoolAttr(false);
 
     if (!isa<LLVM::LLVMArrayType>(operandType)) {
-      LLVM::ConstantOp zero = rewriter.create<LLVM::ConstantOp>(loc, boolZero);
       rewriter.replaceOpWithNewOp<LLVMOp>(op, resultType, adaptor.getOperand(),
-                                          zero);
+                                          false);
       return success();
     }
 
     auto vectorType = dyn_cast<VectorType>(resultType);
     if (!vectorType)
       return failure();
 
     return LLVM::detail::handleMultidimensionalVectors(
         op.getOperation(), adaptor.getOperands(), *this->getTypeConverter(),
         [&](Type llvm1DVectorTy, ValueRange operands) {
-          LLVM::ConstantOp zero =
-              rewriter.create<LLVM::ConstantOp>(loc, boolZero);
           return rewriter.create<LLVMOp>(loc, llvm1DVectorTy, operands[0],
-                                         zero);
+                                         false);
         },
         rewriter);
   }
 };
 
 using CountLeadingZerosOpLowering =
     IntOpWithFlagLowering<math::CountLeadingZerosOp, LLVM::CountLeadingZerosOp>;
 using CountTrailingZerosOpLowering =
-    IntOpWithFlagLowering<math::CountTrailingZerosOp, LLVM::CountTrailingZerosOp>;
+    IntOpWithFlagLowering<math::CountTrailingZerosOp,
+                          LLVM::CountTrailingZerosOp>;
 using AbsIOpLowering = IntOpWithFlagLowering<math::AbsIOp, LLVM::AbsOp>;
 
 // A `expm1` is converted into `exp - 1`.
 struct ExpM1OpLowering : public ConvertOpToLLVMPattern<math::ExpM1Op> {
   using ConvertOpToLLVMPattern<math::ExpM1Op>::ConvertOpToLLVMPattern;
 
   LogicalResult
   matchAndRewrite(math::ExpM1Op op, OpAdaptor adaptor,

mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp

LogicalResult matchAndRewrite(memref::ReallocOp op, OpAdaptor adaptor,
     // type size computed from dst memref.
     auto srcByteSize =
         rewriter.create<LLVM::MulOp>(loc, srcNumElements, sizeInBytes);
     // Allocate a new buffer.
     auto [dstRawPtr, dstAlignedPtr] =
         allocateBuffer(rewriter, loc, dstByteSize, op);
     // Copy the data from the old buffer to the new buffer.
     Value srcAlignedPtr = desc.alignedPtr(rewriter, loc);
-    Value isVolatile =
-        rewriter.create<LLVM::ConstantOp>(loc, rewriter.getBoolAttr(false));
     auto toVoidPtr = [&](Value ptr) -> Value {
       if (getTypeConverter()->useOpaquePointers())
         return ptr;
       return rewriter.create<LLVM::BitcastOp>(loc, getVoidPtrType(), ptr);
     };
     rewriter.create<LLVM::MemcpyOp>(loc, toVoidPtr(dstAlignedPtr),
                                     toVoidPtr(srcAlignedPtr), srcByteSize,
-                                    isVolatile);
+                                    /*isVolatile=*/false);
     // Deallocate the old buffer.
     LLVM::LLVMFuncOp freeFunc =
         getFreeFn(getTypeConverter(), op->getParentOfType<ModuleOp>());
     rewriter.create<LLVM::CallOp>(loc, freeFunc,
                                   toVoidPtr(desc.allocatedPtr(rewriter, loc)));
     // Replace the old buffer addresses in the MemRefDescriptor with the new
     // buffer addresses.
     desc.setAllocatedPtr(rewriter, loc, dstRawPtr);
matchAndRewrite(memref::PrefetchOp prefetchOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto type = prefetchOp.getMemRefType();
     auto loc = prefetchOp.getLoc();
 
     Value dataPtr = getStridedElementPtr(loc, type, adaptor.getMemref(),
                                          adaptor.getIndices(), rewriter);
 
     // Replace with llvm.prefetch.
-    auto llvmI32Type = typeConverter->convertType(rewriter.getIntegerType(32));
-    auto isWrite = rewriter.create<LLVM::ConstantOp>(loc, llvmI32Type,
-                                                     prefetchOp.getIsWrite());
-    auto localityHint = rewriter.create<LLVM::ConstantOp>(
-        loc, llvmI32Type, prefetchOp.getLocalityHint());
-    auto isData = rewriter.create<LLVM::ConstantOp>(
-        loc, llvmI32Type, prefetchOp.getIsDataCache());
-
+    IntegerAttr isWrite = rewriter.getI32IntegerAttr(prefetchOp.getIsWrite());
+    IntegerAttr localityHint = prefetchOp.getLocalityHintAttr();
+    IntegerAttr isData =
+        rewriter.getI32IntegerAttr(prefetchOp.getIsDataCache());
     rewriter.replaceOpWithNewOp<LLVM::Prefetch>(prefetchOp, dataPtr, isWrite,
                                                 localityHint, isData);
     return success();
   }
 };
 
 struct RankOpLowering : public ConvertOpToLLVMPattern<memref::RankOp> {
   using ConvertOpToLLVMPattern<memref::RankOp>::ConvertOpToLLVMPattern;
lowerToMemCopyIntrinsic(memref::CopyOp op, OpAdaptor adaptor,
     Value srcOffset = srcDesc.offset(rewriter, loc);
     Value srcPtr = rewriter.create<LLVM::GEPOp>(
         loc, srcBasePtr.getType(), elementType, srcBasePtr, srcOffset);
     MemRefDescriptor targetDesc(adaptor.getTarget());
     Value targetBasePtr = targetDesc.alignedPtr(rewriter, loc);
     Value targetOffset = targetDesc.offset(rewriter, loc);
     Value targetPtr = rewriter.create<LLVM::GEPOp>(
         loc, targetBasePtr.getType(), elementType, targetBasePtr, targetOffset);
-    Value isVolatile =
-        rewriter.create<LLVM::ConstantOp>(loc, rewriter.getBoolAttr(false));
     rewriter.create<LLVM::MemcpyOp>(loc, targetPtr, srcPtr, totalSize,
-                                    isVolatile);
+                                    /*isVolatile=*/false);
     rewriter.eraseOp(op);
 
     return success();
   }
 
   LogicalResult
   lowerToMemCopyFunctionCall(memref::CopyOp op, OpAdaptor adaptor,
                              ConversionPatternRewriter &rewriter) const {
if (auto resultTypeU = dyn_cast<UnrankedMemRefType>(resultType)) {
 
       int64_t bytesToSkip =
           2 *
           ceilDiv(getTypeConverter()->getPointerBitwidth(resultAddrSpace), 8);
       Value bytesToSkipConst = rewriter.create<LLVM::ConstantOp>(
           loc, getIndexType(), rewriter.getIndexAttr(bytesToSkip));
       Value copySize = rewriter.create<LLVM::SubOp>(
           loc, getIndexType(), resultUnderlyingSize, bytesToSkipConst);
-      Type llvmBool = typeConverter->convertType(rewriter.getI1Type());
-      Value nonVolatile = rewriter.create<LLVM::ConstantOp>(
-          loc, llvmBool, rewriter.getBoolAttr(false));
       rewriter.create<LLVM::MemcpyOp>(loc, resultIndexVals, sourceIndexVals,
-                                      copySize, nonVolatile);
+                                      copySize, /*isVolatile=*/false);
 
       rewriter.replaceOp(op, ValueRange{result});
       return success();
     }
     return rewriter.notifyMatchFailure(loc, "unexpected memref type");
   }
 };
 

mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp

 static unsigned calculateGlobalIndex(spirv::GlobalVariableOp op) {
   IntegerAttr binding = op->getAttrOfType<IntegerAttr>(bindingName());
   return binding.getInt();
 }
 
 /// Copies the given number of bytes from src to dst pointers.
 static void copy(Location loc, Value dst, Value src, Value size,
                  OpBuilder &builder) {
-  MLIRContext *context = builder.getContext();
-  auto llvmI1Type = IntegerType::get(context, 1);
-  Value isVolatile = builder.create<LLVM::ConstantOp>(
-      loc, llvmI1Type, builder.getBoolAttr(false));
-  builder.create<LLVM::MemcpyOp>(loc, dst, src, size, isVolatile);
+  builder.create<LLVM::MemcpyOp>(loc, dst, src, size, /*isVolatile=*/false);
 }
 
 /// Encodes the binding and descriptor set numbers into a new symbolic name.
 /// The name is specified by
 ///   {kernel_module_name}_{variable_name}_descriptor_set{ds}_binding{b}
 /// to avoid symbolic conflicts, where 'ds' and 'b' are descriptor set and
 /// binding numbers.
 static std::string
for (CopyInfo info : copyInfo)
       copy(loc, info.src, info.dst, info.size, rewriter);
     return success();
   }
 };
 
 class LowerHostCodeToLLVM
     : public impl::LowerHostCodeToLLVMPassBase<LowerHostCodeToLLVM> {
 public:
-
   using Base::Base;
 
   void runOnOperation() override {
     ModuleOp module = getOperation();
 
     // Erase the GPU module.
     for (auto gpuModule :
          llvm::make_early_inc_range(module.getOps<gpu::GPUModuleOp>()))

mlir/lib/Dialect/LLVMIR/IR/LLVMInlining.cpp

Value allocaOp;
     Value one = builder.create<LLVM::ConstantOp>(loc, builder.getI64Type(),
                                                  builder.getI64IntegerAttr(1));
     allocaOp = builder.create<LLVM::AllocaOp>(
         loc, argument.getType(), elementType, one, targetAlignment);
   }
   // Copy the pointee to the newly allocated value.
   Value copySize = builder.create<LLVM::ConstantOp>(
       loc, builder.getI64Type(), builder.getI64IntegerAttr(elementTypeSize));
-  Value isVolatile = builder.create<LLVM::ConstantOp>(
-      loc, builder.getI1Type(), builder.getBoolAttr(false));
-  builder.create<LLVM::MemcpyOp>(loc, allocaOp, argument, copySize, isVolatile);
+  builder.create<LLVM::MemcpyOp>(loc, allocaOp, argument, copySize,
+                                 /*isVolatile=*/false);
   return allocaOp;
 }
 
 /// Handles a function argument marked with the byval attribute by introducing a
 /// memcpy or realigning the defining operation, if required either due to the
 /// pointee being writeable in the callee, and/or due to an alignment mismatch.
 /// `requestedAlignment` specifies the alignment set in the "align" argument
 /// attribute (or 1 if no align attribute was set).

mlir/lib/Target/LLVMIR/ModuleTranslation.cpp

for (Block &b : region) {
     }
   }
   assert(blocks.size() == region.getBlocks().size() &&
          "some blocks are not sorted");
 
   return blocks;
 }
 
-llvm::Value *mlir::LLVM::detail::createIntrinsicCall(
+llvm::CallInst *mlir::LLVM::detail::createIntrinsicCall(
     llvm::IRBuilderBase &builder, llvm::Intrinsic::ID intrinsic,
     ArrayRef<llvm::Value *> args, ArrayRef<llvm::Type *> tys) {
   llvm::Module *module = builder.GetInsertBlock()->getModule();
   llvm::Function *fn = llvm::Intrinsic::getDeclaration(module, intrinsic, tys);
   return builder.CreateCall(fn, args);
 }
 
 /// Given a single MLIR operation, create the corresponding LLVM IR operation

mlir/test/Conversion/FuncToLLVM/calling-convention.mlir

func.func @return_var_memref_caller(%arg0: memref<4x3xf32>) {
 
   // CHECK: %[[PTR_SIZE:.*]] = llvm.mlir.constant
   // CHECK: %[[DOUBLE_PTR_SIZE:.*]] = llvm.mul %[[TWO]], %[[PTR_SIZE]]
   // CHECK: %[[RANK:.*]] = llvm.extractvalue %[[CALL_RES]][0] : !llvm.struct<(i64, ptr)>
   // CHECK: %[[DOUBLE_RANK:.*]] = llvm.mul %[[TWO]], %[[RANK]]
   // CHECK: %[[DOUBLE_RANK_INC:.*]] = llvm.add %[[DOUBLE_RANK]], %[[ONE]]
   // CHECK: %[[TABLES_SIZE:.*]] = llvm.mul %[[DOUBLE_RANK_INC]], %[[IDX_SIZE]]
   // CHECK: %[[ALLOC_SIZE:.*]] = llvm.add %[[DOUBLE_PTR_SIZE]], %[[TABLES_SIZE]]
-  // CHECK: %[[FALSE:.*]] = llvm.mlir.constant(false)
   // CHECK: %[[ALLOCA:.*]] = llvm.alloca %[[ALLOC_SIZE]] x i8
   // CHECK: %[[SOURCE:.*]] = llvm.extractvalue %[[CALL_RES]][1]
-  // CHECK: "llvm.intr.memcpy"(%[[ALLOCA]], %[[SOURCE]], %[[ALLOC_SIZE]], %[[FALSE]])
+  // CHECK: "llvm.intr.memcpy"(%[[ALLOCA]], %[[SOURCE]], %[[ALLOC_SIZE]]) <{isVolatile = false}>
   // CHECK: llvm.call @free(%[[SOURCE]])
   // CHECK: %[[DESC:.*]] = llvm.mlir.undef : !llvm.struct<(i64, ptr)>
   // CHECK: %[[RANK:.*]] = llvm.extractvalue %[[CALL_RES]][0] : !llvm.struct<(i64, ptr)>
   // CHECK: %[[DESC_1:.*]] = llvm.insertvalue %[[RANK]], %[[DESC]][0]
   // CHECK: llvm.insertvalue %[[ALLOCA]], %[[DESC_1]][1]
   return
 }
 
func.func @return_var_memref(%arg0: memref<4x3xf32>) -> memref<*xf32> attributes { llvm.emit_c_interface } {
   // CHECK: %[[IDX_SIZE:.*]] = llvm.mlir.constant
 
   // CHECK: %[[PTR_SIZE:.*]] = llvm.mlir.constant
   // CHECK: %[[DOUBLE_PTR_SIZE:.*]] = llvm.mul %[[TWO]], %[[PTR_SIZE]]
   // CHECK: %[[DOUBLE_RANK:.*]] = llvm.mul %[[TWO]], %[[RANK]]
   // CHECK: %[[DOUBLE_RANK_INC:.*]] = llvm.add %[[DOUBLE_RANK]], %[[ONE]]
   // CHECK: %[[TABLES_SIZE:.*]] = llvm.mul %[[DOUBLE_RANK_INC]], %[[IDX_SIZE]]
   // CHECK: %[[ALLOC_SIZE:.*]] = llvm.add %[[DOUBLE_PTR_SIZE]], %[[TABLES_SIZE]]
-  // CHECK: %[[FALSE:.*]] = llvm.mlir.constant(false)
   // CHECK: %[[ALLOCATED:.*]] = llvm.call @malloc(%[[ALLOC_SIZE]])
-  // CHECK: "llvm.intr.memcpy"(%[[ALLOCATED]], %[[ALLOCA]], %[[ALLOC_SIZE]], %[[FALSE]])
+  // CHECK: "llvm.intr.memcpy"(%[[ALLOCATED]], %[[ALLOCA]], %[[ALLOC_SIZE]]) <{isVolatile = false}>
   // CHECK: %[[NEW_DESC:.*]] = llvm.mlir.undef : !llvm.struct<(i64, ptr)>
   // CHECK: %[[NEW_DESC_1:.*]] = llvm.insertvalue %[[RANK]], %[[NEW_DESC]][0]
   // CHECK: %[[NEW_DESC_2:.*]] = llvm.insertvalue %[[ALLOCATED]], %[[NEW_DESC_1]][1]
   // CHECK: llvm.return %[[NEW_DESC_2]]
   return %0 : memref<*xf32>
 }
 
 // Check that the result memref is passed as parameter
 // CHECK-LABEL: @_mlir_ciface_return_var_memref
 // CHECK-SAME: (%{{.*}}: !llvm.ptr, %{{.*}}: !llvm.ptr)
 
 // CHECK-LABEL: llvm.func @return_two_var_memref_caller
 func.func @return_two_var_memref_caller(%arg0: memref<4x3xf32>) {
   // Only check that we create two different descriptors using different
   // memory, and deallocate both sources. The size computation is same as for
   // the single result.
   // CHECK: %[[CALL_RES:.*]] = llvm.call @return_two_var_memref
   // CHECK: %[[RES_1:.*]] = llvm.extractvalue %[[CALL_RES]][0]
   // CHECK: %[[RES_2:.*]] = llvm.extractvalue %[[CALL_RES]][1]
   %0:2 = call @return_two_var_memref(%arg0) : (memref<4x3xf32>) -> (memref<*xf32>, memref<*xf32>)
 
   // CHECK: %[[ALLOCA_1:.*]] = llvm.alloca %{{.*}} x i8
   // CHECK: %[[SOURCE_1:.*]] = llvm.extractvalue %[[RES_1:.*]][1] : ![[DESC_TYPE:.*>]]
-  // CHECK: "llvm.intr.memcpy"(%[[ALLOCA_1]], %[[SOURCE_1]], %{{.*}}, %[[FALSE:.*]])
+  // CHECK: "llvm.intr.memcpy"(%[[ALLOCA_1]], %[[SOURCE_1]], %{{.*}}) <{isVolatile = false}>
   // CHECK: llvm.call @free(%[[SOURCE_1]])
   // CHECK: %[[DESC_1:.*]] = llvm.mlir.undef : ![[DESC_TYPE]]
   // CHECK: %[[DESC_11:.*]] = llvm.insertvalue %{{.*}}, %[[DESC_1]][0]
   // CHECK: llvm.insertvalue %[[ALLOCA_1]], %[[DESC_11]][1]
 
   // CHECK: %[[ALLOCA_2:.*]] = llvm.alloca %{{.*}} x i8
   // CHECK: %[[SOURCE_2:.*]] = llvm.extractvalue %[[RES_2:.*]][1]
-  // CHECK: "llvm.intr.memcpy"(%[[ALLOCA_2]], %[[SOURCE_2]], %{{.*}}, %[[FALSE]])
+  // CHECK: "llvm.intr.memcpy"(%[[ALLOCA_2]], %[[SOURCE_2]], %{{.*}}) <{isVolatile = false}>
   // CHECK: llvm.call @free(%[[SOURCE_2]])
   // CHECK: %[[DESC_2:.*]] = llvm.mlir.undef : ![[DESC_TYPE]]
   // CHECK: %[[DESC_21:.*]] = llvm.insertvalue %{{.*}}, %[[DESC_2]][0]
   // CHECK: llvm.insertvalue %[[ALLOCA_2]], %[[DESC_21]][1]
   return
 }
 
 // CHECK-LABEL: llvm.func @return_two_var_memref
 func.func @return_two_var_memref(%arg0: memref<4x3xf32>) -> (memref<*xf32>, memref<*xf32>) attributes { llvm.emit_c_interface } {
   // Match the construction of the unranked descriptor.
   // CHECK: %[[ALLOCA:.*]] = llvm.alloca
   // CHECK: %[[DESC_0:.*]] = llvm.mlir.undef : !llvm.struct<(i64, ptr)>
   // CHECK: %[[DESC_1:.*]] = llvm.insertvalue %{{.*}}, %[[DESC_0]][0]
   // CHECK: %[[DESC_2:.*]] = llvm.insertvalue %[[ALLOCA]], %[[DESC_1]][1]
   %0 = memref.cast %arg0 : memref<4x3xf32> to memref<*xf32>
 
   // Only check that we allocate the memory for each operand of the "return"
   // separately, even if both operands are the same value. The calling
   // convention requires the caller to free them and the caller cannot know
   // whether they are the same value or not.
   // CHECK: %[[ALLOCATED_1:.*]] = llvm.call @malloc(%{{.*}})
-  // CHECK: "llvm.intr.memcpy"(%[[ALLOCATED_1]], %[[ALLOCA]], %{{.*}}, %[[FALSE:.*]])
+  // CHECK: "llvm.intr.memcpy"(%[[ALLOCATED_1]], %[[ALLOCA]], %{{.*}}) <{isVolatile = false}>
   // CHECK: %[[RES_1:.*]] = llvm.mlir.undef
   // CHECK: %[[RES_11:.*]] = llvm.insertvalue %{{.*}}, %[[RES_1]][0]
   // CHECK: %[[RES_12:.*]] = llvm.insertvalue %[[ALLOCATED_1]], %[[RES_11]][1]
 
   // CHECK: %[[ALLOCATED_2:.*]] = llvm.call @malloc(%{{.*}})
-  // CHECK: "llvm.intr.memcpy"(%[[ALLOCATED_2]], %[[ALLOCA]], %{{.*}}, %[[FALSE]])
+  // CHECK: "llvm.intr.memcpy"(%[[ALLOCATED_2]], %[[ALLOCA]], %{{.*}}) <{isVolatile = false}>
   // CHECK: %[[RES_2:.*]] = llvm.mlir.undef
   // CHECK: %[[RES_21:.*]] = llvm.insertvalue %{{.*}}, %[[RES_2]][0]
   // CHECK: %[[RES_22:.*]] = llvm.insertvalue %[[ALLOCATED_2]], %[[RES_21]][1]
 
   // CHECK: %[[RESULTS:.*]] = llvm.mlir.undef : !llvm.struct<(struct<(i64, ptr)>, struct<(i64, ptr)>)>
   // CHECK: %[[RESULTS_1:.*]] = llvm.insertvalue %[[RES_12]], %[[RESULTS]]
   // CHECK: %[[RESULTS_2:.*]] = llvm.insertvalue %[[RES_22]], %[[RESULTS_1]]
   // CHECK: llvm.return %[[RESULTS_2]]

mlir/test/Conversion/MathToLLVM/math-to-llvm.mlir

func.func @ops(%arg0: f32, %arg1: f32, %arg2: i32, %arg3: i32, %arg4: f64) {
   // CHECK: = llvm.intr.sqrt(%{{.*}}) : (f64) -> f64
   %3 = math.sqrt %arg4 : f64
   func.return
 }
 
 // -----
 
 func.func @absi(%arg0: i32) -> i32 {
-  // CHECK: %[[FALSE:.*]] = llvm.mlir.constant(false
-  // CHECK: = "llvm.intr.abs"(%{{.*}}, %[[FALSE]]) : (i32, i1) -> i32
+  // CHECK: = "llvm.intr.abs"(%{{.*}}) <{is_int_min_poison = false}> : (i32) -> i32
   %0 = math.absi %arg0 : i32
   return %0 : i32
 }
 
 // -----
 
 // CHECK-LABEL: func @log1p(
 // CHECK-SAME: f32
func.func @sine(%arg0 : f32) {
   func.return
 }
 
 // -----
 
 // CHECK-LABEL: func @ctlz(
 // CHECK-SAME: i32
 func.func @ctlz(%arg0 : i32) {
-  // CHECK: %[[ZERO:.+]] = llvm.mlir.constant(false) : i1
-  // CHECK: "llvm.intr.ctlz"(%arg0, %[[ZERO]]) : (i32, i1) -> i32
+  // CHECK: "llvm.intr.ctlz"(%arg0) <{is_zero_poison = false}> : (i32) -> i32
   %0 = math.ctlz %arg0 : i32
   func.return
 }
 
 // -----
 
 // CHECK-LABEL: func @cttz(
 // CHECK-SAME: i32
 func.func @cttz(%arg0 : i32) {
-  // CHECK: %[[ZERO:.+]] = llvm.mlir.constant(false) : i1
-  // CHECK: "llvm.intr.cttz"(%arg0, %[[ZERO]]) : (i32, i1) -> i32
+  // CHECK: "llvm.intr.cttz"(%arg0) <{is_zero_poison = false}> : (i32) -> i32
   %0 = math.cttz %arg0 : i32
   func.return
 }
 
 // -----
 
 // CHECK-LABEL: func @cttz_vec(
 // CHECK-SAME: i32
 func.func @cttz_vec(%arg0 : vector<4xi32>) {
-  // CHECK: %[[ZERO:.+]] = llvm.mlir.constant(false) : i1
-  // CHECK: "llvm.intr.cttz"(%arg0, %[[ZERO]]) : (vector<4xi32>, i1) -> vector<4xi32>
+  // CHECK: "llvm.intr.cttz"(%arg0) <{is_zero_poison = false}> : (vector<4xi32>) -> vector<4xi32>
   %0 = math.cttz %arg0 : vector<4xi32>
   func.return
 }
 
 // -----
 
 // CHECK-LABEL: func @ctpop(
 // CHECK-SAME: i32

mlir/test/Conversion/MemRefToLLVM/convert-dynamic-memref-ops.mlir

 func.func @prefetch(%A : memref<?x?xf32>, %i : index, %j : index) {
 //      CHECK-DAG:  %[[I:.*]] = builtin.unrealized_conversion_cast %[[Iarg]]
 //      CHECK-DAG:  %[[J:.*]] = builtin.unrealized_conversion_cast %[[Jarg]]
 //      CHECK:  %[[ptr:.*]] = llvm.extractvalue %[[ld:.*]][1] : !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)>
 // CHECK-NEXT:  %[[st0:.*]] = llvm.extractvalue %[[ld]][4, 0] : !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)>
 // CHECK-NEXT:  %[[offI:.*]] = llvm.mul %[[I]], %[[st0]] : i64
 // CHECK-NEXT:  %[[off1:.*]] = llvm.add %[[offI]], %[[J]] : i64
 // CHECK-NEXT:  %[[addr:.*]] = llvm.getelementptr %[[ptr]][%[[off1]]] : (!llvm.ptr, i64) -> !llvm.ptr, f32
-// CHECK-NEXT:  [[C1:%.*]] = llvm.mlir.constant(1 : i32) : i32
-// CHECK-NEXT:  [[C3:%.*]] = llvm.mlir.constant(3 : i32) : i32
-// CHECK-NEXT:  [[C1_1:%.*]] = llvm.mlir.constant(1 : i32) : i32
-// CHECK-NEXT:  "llvm.intr.prefetch"(%[[addr]], [[C1]], [[C3]], [[C1_1]]) : (!llvm.ptr, i32, i32, i32) -> ()
+// CHECK-NEXT:  "llvm.intr.prefetch"(%[[addr]]) <{cache = 1 : i32, hint = 3 : i32, rw = 1 : i32}> : (!llvm.ptr) -> ()
   memref.prefetch %A[%i, %j], write, locality<3>, data : memref<?x?xf32>
-// CHECK:  [[C0:%.*]] = llvm.mlir.constant(0 : i32) : i32
-// CHECK:  [[C0_1:%.*]] = llvm.mlir.constant(0 : i32) : i32
-// CHECK:  [[C1_2:%.*]] = llvm.mlir.constant(1 : i32) : i32
-// CHECK:  "llvm.intr.prefetch"(%{{.*}}, [[C0]], [[C0_1]], [[C1_2]]) : (!llvm.ptr, i32, i32, i32) -> ()
+// CHECK:  "llvm.intr.prefetch"(%{{.*}}) <{cache = 1 : i32, hint = 0 : i32, rw = 0 : i32}> : (!llvm.ptr) -> ()
   memref.prefetch %A[%i, %j], read, locality<0>, data : memref<?x?xf32>
-// CHECK:  [[C0_2:%.*]] = llvm.mlir.constant(0 : i32) : i32
-// CHECK:  [[C2:%.*]] = llvm.mlir.constant(2 : i32) : i32
-// CHECK:  [[C0_3:%.*]] = llvm.mlir.constant(0 : i32) : i32
-// CHECK:  "llvm.intr.prefetch"(%{{.*}}, [[C0_2]], [[C2]], [[C0_3]]) : (!llvm.ptr, i32, i32, i32) -> ()
+// CHECK:  "llvm.intr.prefetch"(%{{.*}}) <{cache = 0 : i32, hint = 2 : i32, rw = 0 : i32}> : (!llvm.ptr) -> ()
   memref.prefetch %A[%i, %j], read, locality<2>, instr : memref<?x?xf32>
   return
 }
 
 // -----
 
 // CHECK-LABEL: func @dynamic_store
 // CHECK:         %{{.*}}, %[[Iarg:.*]]: index, %[[Jarg:.*]]: index

 // CHECK: [[SOURCE_ALIGN_GEP:%.*]] = llvm.getelementptr [[SOURCE_DESC]][1]
 // CHECK: [[SOURCE_ALIGN:%.*]] = llvm.load [[SOURCE_ALIGN_GEP]] : !llvm.ptr
 // CHECK: [[RESULT_ALLOC:%.*]] = llvm.addrspacecast [[SOURCE_ALLOC]] : !llvm.ptr to !llvm.ptr<1>
 // CHECK: [[RESULT_ALIGN:%.*]] = llvm.addrspacecast [[SOURCE_ALIGN]] : !llvm.ptr to !llvm.ptr<1>
 // CHECK: llvm.store [[RESULT_ALLOC]], [[RESULT_DESC]] : !llvm.ptr
 // CHECK: [[RESULT_ALIGN_GEP:%.*]] = llvm.getelementptr [[RESULT_DESC]][1]
 // CHECK: llvm.store [[RESULT_ALIGN]], [[RESULT_ALIGN_GEP]] : !llvm.ptr
 
-// Memcpy remaniing values
+// Memcpy remaining values
 
 // CHECK: [[SOURCE_OFFSET_GEP:%.*]] = llvm.getelementptr [[SOURCE_DESC]][2]
 // CHECK: [[RESULT_OFFSET_GEP:%.*]] = llvm.getelementptr [[RESULT_DESC]][2]
 // CHECK: [[SIZEOF_TWO_RESULT_PTRS:%.*]] = llvm.mlir.constant(16 : index) : i64
 // CHECK: [[COPY_SIZE:%.*]] = llvm.sub [[DESC_ALLOC_SIZE]], [[SIZEOF_TWO_RESULT_PTRS]]
-// CHECK: [[FALSE:%.*]] = llvm.mlir.constant(false) : i1
-// CHECK: "llvm.intr.memcpy"([[RESULT_OFFSET_GEP]], [[SOURCE_OFFSET_GEP]], [[COPY_SIZE]], [[FALSE]])
+// CHECK: "llvm.intr.memcpy"([[RESULT_OFFSET_GEP]], [[SOURCE_OFFSET_GEP]], [[COPY_SIZE]]) <{isVolatile = false}>
 
 // -----
 
 // CHECK-LABEL: func @memref_cast_static_to_dynamic
 func.func @memref_cast_static_to_dynamic(%static : memref<10x42xf32>) {
 // CHECK-NOT: llvm.bitcast
   %0 = memref.cast %static : memref<10x42xf32> to memref<?x?xf32>
   return

 // CHECK:           ^bb1:
 // CHECK:           %[[dst_null:.*]] = llvm.mlir.null : !llvm.ptr
 // CHECK:           %[[dst_gep:.*]] = llvm.getelementptr %[[dst_null]][1]
 // CHECK:           %[[dst_es:.*]] = llvm.ptrtoint %[[dst_gep]] : !llvm.ptr to i64
 // CHECK:           %[[dst_size:.*]] = llvm.mul %[[dst_dim]], %[[dst_es]]
 // CHECK:           %[[src_size:.*]] = llvm.mul %[[src_dim]], %[[dst_es]]
 // CHECK:           %[[new_buffer_raw:.*]] = llvm.call @malloc(%[[dst_size]])
 // CHECK:           %[[old_buffer_aligned:.*]] = llvm.extractvalue %[[descriptor]][1]
-// CHECK:           %[[volatile:.*]] = llvm.mlir.constant(false) : i1
-// CHECK:           "llvm.intr.memcpy"(%[[new_buffer_raw]], %[[old_buffer_aligned]], %[[src_size]], %[[volatile]])
+// CHECK:           "llvm.intr.memcpy"(%[[new_buffer_raw]], %[[old_buffer_aligned]], %[[src_size]]) <{isVolatile = false}>
 // CHECK:           %[[old_buffer_unaligned:.*]] = llvm.extractvalue %[[descriptor]][0]
 // CHECK:           llvm.call @free(%[[old_buffer_unaligned]])
 // CHECK:           %[[descriptor_update1:.*]] = llvm.insertvalue %[[new_buffer_raw]], %[[descriptor]][0]
 // CHECK:           %[[descriptor_update2:.*]] = llvm.insertvalue %[[new_buffer_raw]], %[[descriptor_update1]][1]
 // CHECK:           llvm.br ^bb2(%[[descriptor_update2]]
 // CHECK:           ^bb2(%[[descriptor_update3:.*]]: !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>):
 // CHECK:           %[[descriptor_update4:.*]] = llvm.insertvalue %[[dst_dim]], %[[descriptor_update3]][3, 0]
 // CHECK:           %[[descriptor_update5:.*]] = builtin.unrealized_conversion_cast %[[descriptor_update4]]

 // CHECK:           %[[new_buffer_int:.*]] = llvm.ptrtoint %[[new_buffer_raw]] : !llvm.ptr
 // CHECK:           %[[const_1:.*]] = llvm.mlir.constant(1 : index) : i64
 // CHECK:           %[[alignment_m1:.*]] = llvm.sub %[[alignment]], %[[const_1]]
 // CHECK:           %[[ptr_alignment_m1:.*]] = llvm.add %[[new_buffer_int]], %[[alignment_m1]]
 // CHECK:           %[[padding:.*]] = llvm.urem %[[ptr_alignment_m1]], %[[alignment]]
 // CHECK:           %[[new_buffer_aligned_int:.*]] = llvm.sub %[[ptr_alignment_m1]], %[[padding]]
 // CHECK:           %[[new_buffer_aligned:.*]] = llvm.inttoptr %[[new_buffer_aligned_int]] : i64 to !llvm.ptr
 // CHECK:           %[[old_buffer_aligned:.*]] = llvm.extractvalue %[[descriptor]][1]
-// CHECK:           %[[volatile:.*]] = llvm.mlir.constant(false) : i1
-// CHECK:           "llvm.intr.memcpy"(%[[new_buffer_aligned]], %[[old_buffer_aligned]], %[[src_size]], %[[volatile]])
+// CHECK:           "llvm.intr.memcpy"(%[[new_buffer_aligned]], %[[old_buffer_aligned]], %[[src_size]]) <{isVolatile = false}>
 // CHECK:           %[[old_buffer_unaligned:.*]] = llvm.extractvalue %[[descriptor]][0]
 // CHECK:           llvm.call @free(%[[old_buffer_unaligned]])
 // CHECK:           %[[descriptor_update1:.*]] = llvm.insertvalue %[[new_buffer_raw]], %[[descriptor]][0]
 // CHECK:           %[[descriptor_update2:.*]] = llvm.insertvalue %[[new_buffer_aligned]], %[[descriptor_update1]][1]
 // CHECK:           llvm.br ^bb2(%[[descriptor_update2]]
 // CHECK:           ^bb2(%[[descriptor_update3:.*]]: !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>):
 // CHECK:           %[[descriptor_update4:.*]] = llvm.insertvalue %[[dst_dim]], %[[descriptor_update3]][3, 0]
 // CHECK:           %[[descriptor_update5:.*]] = builtin.unrealized_conversion_cast %[[descriptor_update4]]

 // ALIGNED-ALLOC-DAG:       %[[alignment:.*]] = llvm.mlir.constant(8 : index) : i64
 // ALIGNED-ALLOC-DAG:       %[[const_1:.*]] = llvm.mlir.constant(1 : index) : i64
 // ALIGNED-ALLOC:           %[[alignment_m1:.*]] = llvm.sub %[[alignment]], %[[const_1]]
 // ALIGNED-ALLOC:           %[[size_alignment_m1:.*]] = llvm.add %[[dst_size]], %[[alignment_m1]]
 // ALIGNED-ALLOC:           %[[padding:.*]] = llvm.urem %[[size_alignment_m1]], %[[alignment]]
 // ALIGNED-ALLOC:           %[[adjust_dst_size:.*]] = llvm.sub %[[size_alignment_m1]], %[[padding]]
 // ALIGNED-ALLOC:           %[[new_buffer_raw:.*]] = llvm.call @aligned_alloc(%[[alignment]], %[[adjust_dst_size]])
 // ALIGNED-ALLOC:           %[[old_buffer_aligned:.*]] = llvm.extractvalue %[[descriptor]][1]
-// ALIGNED-ALLOC:           %[[volatile:.*]] = llvm.mlir.constant(false) : i1
-// ALIGNED-ALLOC:           "llvm.intr.memcpy"(%[[new_buffer_raw]], %[[old_buffer_aligned]], %[[src_size]], %[[volatile]])
+// ALIGNED-ALLOC:           "llvm.intr.memcpy"(%[[new_buffer_raw]], %[[old_buffer_aligned]], %[[src_size]]) <{isVolatile = false}>
 // ALIGNED-ALLOC:           %[[old_buffer_unaligned:.*]] = llvm.extractvalue %[[descriptor]][0]
 // ALIGNED-ALLOC:           llvm.call @free(%[[old_buffer_unaligned]])
 // ALIGNED-ALLOC:           %[[descriptor_update1:.*]] = llvm.insertvalue %[[new_buffer_raw]], %[[descriptor]][0]
 // ALIGNED-ALLOC:           %[[descriptor_update2:.*]] = llvm.insertvalue %[[new_buffer_raw]], %[[descriptor_update1]][1]
 // ALIGNED-ALLOC:           llvm.br ^bb2(%[[descriptor_update2]] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>)
 // ALIGNED-ALLOC:           ^bb2(%[[descriptor_update3:.*]]: !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>):
 // ALIGNED-ALLOC:           %[[descriptor_update4:.*]] = llvm.insertvalue %[[dst_dim]], %[[descriptor_update3]][3, 0]
 // ALIGNED-ALLOC:           %[[descriptor_update5:.*]] = builtin.unrealized_conversion_cast %[[descriptor_update4]]
 // ALIGNED-ALLOC:           return %[[descriptor_update5]] : memref<?xf32>
 
   %out = memref.realloc %in(%d)  {alignment = 8} : memref<?xf32> to memref<?xf32>
   return %out : memref<?xf32>
 }
 

mlir/test/Conversion/MemRefToLLVM/convert-static-memref-ops.mlir

 // CHECK:           ^bb1:
 // CHECK:           %[[dst_null:.*]] = llvm.mlir.null : !llvm.ptr
 // CHECK:           %[[dst_gep:.*]] = llvm.getelementptr %[[dst_null]][1]
 // CHECK:           %[[dst_es:.*]] = llvm.ptrtoint %[[dst_gep]] : !llvm.ptr to i64
 // CHECK:           %[[dst_size:.*]] = llvm.mul %[[dst_dim]], %[[dst_es]]
 // CHECK:           %[[src_size:.*]] = llvm.mul %[[src_dim]], %[[dst_es]]
 // CHECK:           %[[new_buffer_raw:.*]] = llvm.call @malloc(%[[dst_size]])
 // CHECK:           %[[old_buffer_aligned:.*]] = llvm.extractvalue %[[descriptor]][1]
-// CHECK:           %[[volatile:.*]] = llvm.mlir.constant(false) : i1
-// CHECK:           "llvm.intr.memcpy"(%[[new_buffer_raw]], %[[old_buffer_aligned]], %[[src_size]], %[[volatile]])
+// CHECK:           "llvm.intr.memcpy"(%[[new_buffer_raw]], %[[old_buffer_aligned]], %[[src_size]]) <{isVolatile = false}>
 // CHECK:           %[[old_buffer_unaligned:.*]] = llvm.extractvalue %[[descriptor]][0]
 // CHECK:           llvm.call @free(%[[old_buffer_unaligned]])
 // CHECK:           %[[descriptor_update1:.*]] = llvm.insertvalue %[[new_buffer_raw]], %[[descriptor]][0]
 // CHECK:           %[[descriptor_update2:.*]] = llvm.insertvalue %[[new_buffer_raw]], %[[descriptor_update1]][1]
 // CHECK:           llvm.br ^bb2(%[[descriptor_update2]] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>)
 // CHECK:           ^bb2(%[[descriptor_update3:.*]]: !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>):
 // CHECK:           %[[descriptor_update4:.*]] = llvm.insertvalue %[[dst_dim]], %[[descriptor_update3]][3, 0]
 // CHECK:           %[[descriptor_update5:.*]] = builtin.unrealized_conversion_cast %[[descriptor_update4]]

 // CHECK:           %[[new_buffer_int:.*]] = llvm.ptrtoint %[[new_buffer_raw]] : !llvm.ptr
 // CHECK:           %[[const_1:.*]] = llvm.mlir.constant(1 : index) : i64
 // CHECK:           %[[alignment_m1:.*]] = llvm.sub %[[alignment]], %[[const_1]]
 // CHECK:           %[[ptr_alignment_m1:.*]] = llvm.add %[[new_buffer_int]], %[[alignment_m1]]
 // CHECK:           %[[padding:.*]] = llvm.urem %[[ptr_alignment_m1]], %[[alignment]]
 // CHECK:           %[[new_buffer_aligned_int:.*]] = llvm.sub %[[ptr_alignment_m1]], %[[padding]]
 // CHECK:           %[[new_buffer_aligned:.*]] = llvm.inttoptr %[[new_buffer_aligned_int]] : i64 to !llvm.ptr
 // CHECK:           %[[old_buffer_aligned:.*]] = llvm.extractvalue %[[descriptor]][1]
-// CHECK:           %[[volatile:.*]] = llvm.mlir.constant(false) : i1
-// CHECK:           "llvm.intr.memcpy"(%[[new_buffer_aligned]], %[[old_buffer_aligned]], %[[src_size]], %[[volatile]])
+// CHECK:           "llvm.intr.memcpy"(%[[new_buffer_aligned]], %[[old_buffer_aligned]], %[[src_size]]) <{isVolatile = false}>
 // CHECK:           %[[old_buffer_unaligned:.*]] = llvm.extractvalue %[[descriptor]][0]
 // CHECK:           llvm.call @free(%[[old_buffer_unaligned]])
 // CHECK:           %[[descriptor_update1:.*]] = llvm.insertvalue %[[new_buffer_raw]], %[[descriptor]][0]
 // CHECK:           %[[descriptor_update2:.*]] = llvm.insertvalue %[[new_buffer_aligned]], %[[descriptor_update1]][1]
 // CHECK:           llvm.br ^bb2(%[[descriptor_update2]]
 // CHECK:           ^bb2(%[[descriptor_update3:.*]]: !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>):
 // CHECK:           %[[descriptor_update4:.*]] = llvm.insertvalue %[[dst_dim]], %[[descriptor_update3]][3, 0]
 // CHECK:           %[[descriptor_update5:.*]] = builtin.unrealized_conversion_cast %[[descriptor_update4]]

mlir/test/Conversion/MemRefToLLVM/memref-to-llvm.mlir

func.func @memref_copy_ranked() {
   // CHECK: [[PTRTOINT:%.*]] = llvm.ptrtoint [[GEP]] : !llvm.ptr to i64
   // CHECK: [[SIZE:%.*]] = llvm.mul [[MUL]], [[PTRTOINT]] : i64
   // CHECK: [[EXTRACT1P:%.*]] = llvm.extractvalue {{%.*}}[1] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
   // CHECK: [[EXTRACT1O:%.*]] = llvm.extractvalue {{%.*}}[2] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
   // CHECK: [[GEP1:%.*]] = llvm.getelementptr [[EXTRACT1P]][[[EXTRACT1O]]] : (!llvm.ptr, i64) -> !llvm.ptr, f32
   // CHECK: [[EXTRACT2P:%.*]] = llvm.extractvalue {{%.*}}[1] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
   // CHECK: [[EXTRACT2O:%.*]] = llvm.extractvalue {{%.*}}[2] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
   // CHECK: [[GEP2:%.*]] = llvm.getelementptr [[EXTRACT2P]][[[EXTRACT2O]]] : (!llvm.ptr, i64) -> !llvm.ptr, f32
-  // CHECK: [[VOLATILE:%.*]] = llvm.mlir.constant(false) : i1
-  // CHECK: "llvm.intr.memcpy"([[GEP2]], [[GEP1]], [[SIZE]], [[VOLATILE]]) : (!llvm.ptr, !llvm.ptr, i64, i1) -> ()
+  // CHECK: "llvm.intr.memcpy"([[GEP2]], [[GEP1]], [[SIZE]]) <{isVolatile = false}>
   return
 }
 
 
 // -----
 
 // CHECK-LABEL: func @memref_copy_contiguous
 func.func @memref_copy_contiguous(%in: memref<16x4xi32>, %offset: index) {
func.func @memref_copy_contiguous(%in: memref<16x4xi32>, %offset: index) {
   // CHECK: [[PTRTOINT:%.*]] = llvm.ptrtoint [[GEP]] : !llvm.ptr to i64
   // CHECK: [[SIZE:%.*]] = llvm.mul [[MUL2]], [[PTRTOINT]] : i64
   // CHECK: [[EXTRACT1P:%.*]] = llvm.extractvalue {{%.*}}[1] : !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)>
   // CHECK: [[EXTRACT1O:%.*]] = llvm.extractvalue {{%.*}}[2] : !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)>
   // CHECK: [[GEP1:%.*]] = llvm.getelementptr [[EXTRACT1P]][[[EXTRACT1O]]] : (!llvm.ptr, i64) -> !llvm.ptr, i32
   // CHECK: [[EXTRACT2P:%.*]] = llvm.extractvalue {{%.*}}[1] : !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)>
   // CHECK: [[EXTRACT2O:%.*]] = llvm.extractvalue {{%.*}}[2] : !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)>
   // CHECK: [[GEP2:%.*]] = llvm.getelementptr [[EXTRACT2P]][[[EXTRACT2O]]] : (!llvm.ptr, i64) -> !llvm.ptr, i32
-  // CHECK: [[VOLATILE:%.*]] = llvm.mlir.constant(false) : i1
-  // CHECK: "llvm.intr.memcpy"([[GEP2]], [[GEP1]], [[SIZE]], [[VOLATILE]]) : (!llvm.ptr, !llvm.ptr, i64, i1) -> ()
+  // CHECK: "llvm.intr.memcpy"([[GEP2]], [[GEP1]], [[SIZE]]) <{isVolatile = false}>
   return
 }
 
 // -----
 
 // CHECK-LABEL: func @memref_copy_0d_offset
 func.func @memref_copy_0d_offset(%in: memref<2xi32>) {
   %buf = memref.alloc() : memref<i32>

mlir/test/Conversion/MemRefToLLVM/typed-pointers.mlir

 // CHECK:           %[[dst_null:.*]] = llvm.mlir.null : !llvm.ptr<f32>
 // CHECK:           %[[dst_gep:.*]] = llvm.getelementptr %[[dst_null]][1]
 // CHECK:           %[[dst_es:.*]] = llvm.ptrtoint %[[dst_gep]] : !llvm.ptr<f32> to i64
 // CHECK:           %[[dst_size:.*]] = llvm.mul %[[dst_dim]], %[[dst_es]]
 // CHECK:           %[[src_size:.*]] = llvm.mul %[[src_dim]], %[[dst_es]]
 // CHECK:           %[[new_buffer_raw:.*]] = llvm.call @malloc(%[[dst_size]])
 // CHECK:           %[[new_buffer:.*]] = llvm.bitcast %[[new_buffer_raw]] : !llvm.ptr<i8> to !llvm.ptr<f32>
 // CHECK:           %[[old_buffer_aligned:.*]] = llvm.extractvalue %[[descriptor]][1]
-// CHECK:           %[[volatile:.*]] = llvm.mlir.constant(false) : i1
 // CHECK-DAG:       %[[new_buffer_void:.*]] = llvm.bitcast %[[new_buffer]] : !llvm.ptr<f32> to !llvm.ptr<i8>
 // CHECK-DAG:       %[[old_buffer_void:.*]] = llvm.bitcast %[[old_buffer_aligned]] : !llvm.ptr<f32> to !llvm.ptr<i8>
-// CHECK:           "llvm.intr.memcpy"(%[[new_buffer_void]], %[[old_buffer_void]], %[[src_size]], %[[volatile]])
+// CHECK:           "llvm.intr.memcpy"(%[[new_buffer_void]], %[[old_buffer_void]], %[[src_size]]) <{isVolatile = false}>
 // CHECK:           %[[old_buffer_unaligned:.*]] = llvm.extractvalue %[[descriptor]][0]
 // CHECK:           %[[old_buffer_unaligned_void:.*]] = llvm.bitcast %[[old_buffer_unaligned]] : !llvm.ptr<f32> to !llvm.ptr<i8>
 // CHECK:           llvm.call @free(%[[old_buffer_unaligned_void]])
 // CHECK:           %[[descriptor_update1:.*]] = llvm.insertvalue %[[new_buffer]], %[[descriptor]][0]
 // CHECK:           %[[descriptor_update2:.*]] = llvm.insertvalue %[[new_buffer]], %[[descriptor_update1]][1]
 // CHECK:           llvm.br ^bb2(%[[descriptor_update2]]
 // CHECK:           ^bb2(%[[descriptor_update3:.*]]: !llvm.struct<(ptr<f32>, ptr<f32>, i64, array<1 x i64>, array<1 x i64>)>):
 // CHECK:           %[[descriptor_update4:.*]] = llvm.insertvalue %[[dst_dim]], %[[descriptor_update3]][3, 0]
 // CHECK:           %[[descriptor_update5:.*]] = builtin.unrealized_conversion_cast %[[descriptor_update4]]
 // CHECK:           return %[[descriptor_update5]] : memref<?xf32>
 
   %out = memref.realloc %in(%d) : memref<?xf32> to memref<?xf32>
   return %out : memref<?xf32>
 }
 

mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir

module attributes {gpu.container_module, spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]>, #spirv.resource_limits<max_compute_workgroup_invocations = 128, max_compute_workgroup_size = [128, 128, 64]>>} {
   //       CHECK:   spirv.func @__spv__foo_bar
 
   //       CHECK:   spirv.EntryPoint "GLCompute" @__spv__foo_bar
   //       CHECK:   spirv.ExecutionMode @__spv__foo_bar "LocalSize", 1, 1, 1
 
   // CHECK-LABEL: @main
   //       CHECK:   %[[SRC:.*]] = llvm.extractvalue %{{.*}}[0] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
   //  CHECK-NEXT:   %[[DEST:.*]] = llvm.mlir.addressof @__spv__foo_bar_arg_0_descriptor_set0_binding0 : !llvm.ptr
-  //  CHECK-NEXT:   llvm.mlir.constant(false) : i1
-  //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[DEST]], %[[SRC]], %[[SIZE:.*]], %{{.*}}) : (!llvm.ptr, !llvm.ptr, i64, i1) -> ()
+  //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[DEST]], %[[SRC]], %[[SIZE:.*]]) <{isVolatile = false}> : (!llvm.ptr, !llvm.ptr, i64) -> ()
   //  CHECK-NEXT:   llvm.call @__spv__foo_bar() : () -> ()
-  //  CHECK-NEXT:   llvm.mlir.constant(false) : i1
-  //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[SRC]], %[[DEST]], %[[SIZE]], %{{.*}}) : (!llvm.ptr, !llvm.ptr, i64, i1) -> ()
+  //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[SRC]], %[[DEST]], %[[SIZE]]) <{isVolatile = false}> : (!llvm.ptr, !llvm.ptr, i64) -> ()
 
   spirv.module @__spv__foo Logical GLSL450 requires #spirv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]> {
     spirv.GlobalVariable @bar_arg_0 bind(0, 0) : !spirv.ptr<!spirv.struct<(!spirv.array<6 x i32, stride=4> [0])>, StorageBuffer>
     spirv.func @bar() "None" attributes {workgroup_attributions = 0 : i64} {
       %0 = spirv.mlir.addressof @bar_arg_0 : !spirv.ptr<!spirv.struct<(!spirv.array<6 x i32, stride=4> [0])>, StorageBuffer>
       spirv.Return
     }
     spirv.EntryPoint "GLCompute" @bar

mlir/test/Dialect/LLVMIR/inlining.mlir

 // RUN: mlir-opt %s -inline -split-input-file | FileCheck %s
 
 #file = #llvm.di_file<"foo.mlir" in "/foo/">
 #variable = #llvm.di_local_variable<scope = #file>
 #variableAddr = #llvm.di_local_variable<scope = #file>
 
 func.func @inner_func_inlinable(%ptr : !llvm.ptr) -> i32 {
   %0 = llvm.mlir.constant(42 : i32) : i32
   %stack = llvm.intr.stacksave : !llvm.ptr
   llvm.store %0, %ptr { alignment = 8 } : i32, !llvm.ptr
   %1 = llvm.load %ptr { alignment = 8 } : !llvm.ptr -> i32
   llvm.intr.dbg.value #variable = %0 : i32
   llvm.intr.dbg.declare #variableAddr = %ptr : !llvm.ptr
   %byte = llvm.mlir.constant(43 : i8) : i8
   %true = llvm.mlir.constant(1 : i1) : i1
-  "llvm.intr.memset"(%ptr, %byte, %0, %true) : (!llvm.ptr, i8, i32, i1) -> ()
-  "llvm.intr.memmove"(%ptr, %ptr, %0, %true) : (!llvm.ptr, !llvm.ptr, i32, i1) -> ()
-  "llvm.intr.memcpy"(%ptr, %ptr, %0, %true) : (!llvm.ptr, !llvm.ptr, i32, i1) -> ()
+  "llvm.intr.memset"(%ptr, %byte, %0) <{isVolatile = true}> : (!llvm.ptr, i8, i32) -> ()
+  "llvm.intr.memmove"(%ptr, %ptr, %0) <{isVolatile = true}> : (!llvm.ptr, !llvm.ptr, i32) -> ()
+  "llvm.intr.memcpy"(%ptr, %ptr, %0) <{isVolatile = true}> : (!llvm.ptr, !llvm.ptr, i32) -> ()
   "llvm.intr.assume"(%true) : (i1) -> ()
   llvm.fence release
   %2 = llvm.atomicrmw add %ptr, %0 monotonic : !llvm.ptr, i32
   %3 = llvm.cmpxchg %ptr, %0, %1 acq_rel monotonic : !llvm.ptr, i32
   llvm.inline_asm has_side_effects "foo", "bar" : () -> ()
   llvm.cond_br %true, ^bb1, ^bb2
 ^bb1:
   llvm.unreachable

mlir/test/Dialect/LLVMIR/roundtrip.mlir

// CHECK: llvm.intr.bitreverse(%{{.*}}) : (i32) -> i32
   %32 = llvm.intr.bitreverse(%arg0) : (i32) -> i32
 
 // CHECK: llvm.intr.ctpop(%{{.*}}) : (i32) -> i32
   %33 = llvm.intr.ctpop(%arg0) : (i32) -> i32
 
 // CHECK: llvm.intr.round(%[[FLOAT]]) : (f32) -> f32
   %34 = llvm.intr.round(%arg1) : (f32) -> f32
 
-// CHECK: "llvm.intr.memcpy"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) : (!llvm.ptr, !llvm.ptr, i32, i1) -> ()
-  "llvm.intr.memcpy"(%arg2, %arg3, %arg0, %arg4) : (!llvm.ptr, !llvm.ptr, i32, i1) -> ()
+// CHECK: "llvm.intr.memcpy"(%{{.*}}, %{{.*}}, %{{.*}}) <{isVolatile = false}> : (!llvm.ptr, !llvm.ptr, i32) -> ()
+  "llvm.intr.memcpy"(%arg2, %arg3, %arg0) <{isVolatile = false}> : (!llvm.ptr, !llvm.ptr, i32) -> ()
 
-// CHECK: "llvm.intr.memcpy"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) : (!llvm.ptr, !llvm.ptr, i32, i1) -> ()
-  "llvm.intr.memcpy"(%arg2, %arg3, %arg0, %arg4) : (!llvm.ptr, !llvm.ptr, i32, i1) -> ()
+// CHECK: "llvm.intr.memcpy"(%{{.*}}, %{{.*}}, %{{.*}}) <{isVolatile = false}> : (!llvm.ptr, !llvm.ptr, i32) -> ()
+  "llvm.intr.memcpy"(%arg2, %arg3, %arg0) <{isVolatile = false}> : (!llvm.ptr, !llvm.ptr, i32) -> ()
 
-// CHECK: %[[SZ:.*]] = llvm.mlir.constant
-  %sz = llvm.mlir.constant(10: i64) : i64
-// CHECK: "llvm.intr.memcpy.inline"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) : (!llvm.ptr, !llvm.ptr, i64, i1) -> ()
-  "llvm.intr.memcpy.inline"(%arg2, %arg3, %sz, %arg4) : (!llvm.ptr, !llvm.ptr, i64, i1) -> ()
+// CHECK: "llvm.intr.memcpy.inline"(%{{.*}}, %{{.*}}) <{isVolatile = false, len = 10 : i64}> : (!llvm.ptr, !llvm.ptr) -> ()
+  "llvm.intr.memcpy.inline"(%arg2, %arg3) <{isVolatile = false, len = 10 : i64}> : (!llvm.ptr, !llvm.ptr) -> ()
 
 // CHECK:  llvm.return
   llvm.return
 }
 
 // CHECK-LABEL: @gep
 llvm.func @gep(%ptr: !llvm.ptr, %idx: i64, %ptr2: !llvm.ptr) {
   // CHECK: llvm.getelementptr %{{.*}}[%{{.*}}, 1, 0] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(i32, struct<(i32, f32)>)>

mlir/test/Target/LLVMIR/Import/intrinsic.ll

 ; RUN: mlir-translate -import-llvm %s | FileCheck %s
 
 ; CHECK-LABEL:  llvm.func @fmuladd_test
 define void @fmuladd_test(float %0, float %1, <8 x float> %2, ptr %3) {
   ; CHECK: llvm.intr.fmuladd(%{{.*}}, %{{.*}}, %{{.*}}) : (f32, f32, f32) -> f32
   %5 = call float @llvm.fmuladd.f32(float %0, float %1, float %0)
   ; CHECK: llvm.intr.fmuladd(%{{.*}}, %{{.*}}, %{{.*}}) : (vector<8xf32>, vector<8xf32>, vector<8xf32>) -> vector<8xf32>
   %6 = call <8 x float> @llvm.fmuladd.v8f32(<8 x float> %2, <8 x float> %2, <8 x float> %2)
   ; CHECK: llvm.intr.fma(%{{.*}}, %{{.*}}, %{{.*}}) : (f32, f32, f32) -> f32
   %7 = call float @llvm.fma.f32(float %0, float %1, float %0)
   ; CHECK: llvm.intr.fma(%{{.*}}, %{{.*}}, %{{.*}}) : (vector<8xf32>, vector<8xf32>, vector<8xf32>) -> vector<8xf32>
   %8 = call <8 x float> @llvm.fma.v8f32(<8 x float> %2, <8 x float> %2, <8 x float> %2)
-  ; CHECK: "llvm.intr.prefetch"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) : (!llvm.ptr, i32, i32, i32) -> ()
+  ; CHECK: "llvm.intr.prefetch"(%{{.*}}) <{cache = 1 : i32, hint = 3 : i32, rw = 0 : i32}> : (!llvm.ptr) -> ()
   call void @llvm.prefetch.p0(ptr %3, i32 0, i32 3, i32 1)
   ret void
 }
 
+; CHECK-LABEL:  llvm.func @fpclass_test
+define void @fpclass_test(float %0, <8 x float> %1) {
+  ; CHECK: "llvm.intr.is.fpclass"(%{{.*}}) <{bit = 0 : i32}> : (f32) -> i1
+  %3 = call i1 @llvm.is.fpclass.f32(float %0, i32 0)
+  ; CHECK: "llvm.intr.is.fpclass"(%{{.*}}) <{bit = 1 : i32}> : (vector<8xf32>) -> vector<8xi1>
+  %4 = call <8 x i1> @llvm.is.fpclass.v8f32(<8 x float> %1, i32 1)
+  ret void
+}
+
 ; CHECK-LABEL:  llvm.func @exp_test
 define void @exp_test(float %0, <8 x float> %1) {
   ; CHECK: llvm.intr.exp(%{{.*}}) : (f32) -> f32
   %3 = call float @llvm.exp.f32(float %0)
   ; CHECK: llvm.intr.exp(%{{.*}}) : (vector<8xf32>) -> vector<8xf32>
   %4 = call <8 x float> @llvm.exp.v8f32(<8 x float> %1)
   ret void
 }
define void @byteswap_test(i32 %0, <8 x i32> %1) {
   ; CHECK:   llvm.intr.bswap(%{{.*}}) : (i32) -> i32
   %3 = call i32 @llvm.bswap.i32(i32 %0)
   ; CHECK:   llvm.intr.bswap(%{{.*}}) : (vector<8xi32>) -> vector<8xi32>
   %4 = call <8 x i32> @llvm.bswap.v8i32(<8 x i32> %1)
   ret void
 }
 ; CHECK-LABEL:  llvm.func @ctlz_test
 define void @ctlz_test(i32 %0, <8 x i32> %1) {
-  ; CHECK:   %[[FALSE:.+]] = llvm.mlir.constant(false) : i1
-  ; CHECK:   "llvm.intr.ctlz"(%{{.*}}, %[[FALSE]]) : (i32, i1) -> i32
+  ; CHECK:   "llvm.intr.ctlz"(%{{.*}}) <{is_zero_poison = false}> : (i32) -> i32
   %3 = call i32 @llvm.ctlz.i32(i32 %0, i1 false)
-  ; CHECK:   "llvm.intr.ctlz"(%{{.*}}, %[[FALSE]]) : (vector<8xi32>, i1) -> vector<8xi32>
+  ; CHECK:   "llvm.intr.ctlz"(%{{.*}}) <{is_zero_poison = false}> : (vector<8xi32>) -> vector<8xi32>
   %4 = call <8 x i32> @llvm.ctlz.v8i32(<8 x i32> %1, i1 false)
   ret void
 }
 ; CHECK-LABEL:  llvm.func @cttz_test
 define void @cttz_test(i32 %0, <8 x i32> %1) {
-  ; CHECK:   %[[FALSE:.+]] = llvm.mlir.constant(false) : i1
-  ; CHECK:   "llvm.intr.cttz"(%{{.*}}, %[[FALSE]]) : (i32, i1) -> i32
+  ; CHECK:   "llvm.intr.cttz"(%{{.*}}) <{is_zero_poison = false}> : (i32) -> i32
   %3 = call i32 @llvm.cttz.i32(i32 %0, i1 false)
-  ; CHECK:   "llvm.intr.cttz"(%{{.*}}, %[[FALSE]]) : (vector<8xi32>, i1) -> vector<8xi32>
+  ; CHECK:   "llvm.intr.cttz"(%{{.*}}) <{is_zero_poison = false}> : (vector<8xi32>) -> vector<8xi32>
   %4 = call <8 x i32> @llvm.cttz.v8i32(<8 x i32> %1, i1 false)
   ret void
 }
+; CHECK-LABEL:  llvm.func @abs_test
+define void @abs_test(i32 %0, <8 x i32> %1) {
+  ; CHECK:   "llvm.intr.abs"(%{{.*}}) <{is_int_min_poison = false}> : (i32) -> i32
+  %3 = call i32 @llvm.abs.i32(i32 %0, i1 false)
+  ; CHECK:   "llvm.intr.abs"(%{{.*}}) <{is_int_min_poison = true}> : (vector<8xi32>) -> vector<8xi32>
+  %4 = call <8 x i32> @llvm.abs.v8i32(<8 x i32> %1, i1 true)
+  ret void
+}
 
 ; CHECK-LABEL:  llvm.func @ctpop_test
 define void @ctpop_test(i32 %0, <8 x i32> %1) {
   ; CHECK:   llvm.intr.ctpop(%{{.*}}) : (i32) -> i32
   %3 = call i32 @llvm.ctpop.i32(i32 %0)
   ; CHECK:   llvm.intr.ctpop(%{{.*}}) : (vector<8xi32>) -> vector<8xi32>
   %4 = call <8 x i32> @llvm.ctpop.v8i32(<8 x i32> %1)
   ret void
define void @trap_intrinsics() {
   ; CHECK: "llvm.intr.ubsantrap"()
   ; CHECK-SAME: failureKind = 1
   call void @llvm.ubsantrap(i8 1)
   ret void
 }
 
 ; CHECK-LABEL:  llvm.func @memcpy_test
 define void @memcpy_test(i32 %0, ptr %1, ptr %2) {
-  ; CHECK: %[[FALSE:.+]] = llvm.mlir.constant(false) : i1
   ; CHECK: %[[CST:.+]] = llvm.mlir.constant(10 : i64) : i64
-  ; CHECK: "llvm.intr.memcpy"(%{{.*}}, %{{.*}}, %{{.*}}, %[[FALSE]]) : (!llvm.ptr, !llvm.ptr, i32, i1) -> ()
+  ; CHECK: "llvm.intr.memcpy"(%{{.*}}, %{{.*}}, %{{.*}}) <{isVolatile = false}> : (!llvm.ptr, !llvm.ptr, i32) -> ()
   call void @llvm.memcpy.p0.p0.i32(ptr %1, ptr %2, i32 %0, i1 false)
-  ; CHECK: "llvm.intr.memcpy.inline"(%{{.*}}, %{{.*}}, %[[CST]], %[[FALSE]]) : (!llvm.ptr, !llvm.ptr, i64, i1) -> ()
+  ; CHECK: "llvm.intr.memcpy.inline"(%{{.*}}, %{{.*}}) <{isVolatile = false, len = 10 : i64}> : (!llvm.ptr, !llvm.ptr) -> ()
   call void @llvm.memcpy.inline.p0.p0.i64(ptr %1, ptr %2, i64 10, i1 false)
+  ; CHECK: "llvm.intr.memcpy.inline"(%{{.*}}, %{{.*}}) <{isVolatile = false, len = 10 : i32}> : (!llvm.ptr, !llvm.ptr) -> ()
+  call void @llvm.memcpy.inline.p0.p0.i32(ptr %1, ptr %2, i32 10, i1 false)
   ret void
 }
 
 ; CHECK-LABEL:  llvm.func @memmove_test
 define void @memmove_test(i32 %0, ptr %1, ptr %2) {
-  ; CHECK: %[[falseval:.+]] = llvm.mlir.constant(false) : i1
-  ; CHECK: "llvm.intr.memmove"(%{{.*}}, %{{.*}}, %{{.*}}, %[[falseval]]) : (!llvm.ptr, !llvm.ptr, i32, i1) -> ()
+  ; CHECK: "llvm.intr.memmove"(%{{.*}}, %{{.*}}, %{{.*}}) <{isVolatile = false}> : (!llvm.ptr, !llvm.ptr, i32) -> ()
   call void @llvm.memmove.p0.p0.i32(ptr %1, ptr %2, i32 %0, i1 false)
   ret void
 }
 
 ; CHECK-LABEL:  llvm.func @memset_test
 define void @memset_test(i32 %0, ptr %1, i8 %2) {
-  ; CHECK: %[[falseval:.+]] = llvm.mlir.constant(false) : i1
-  ; CHECK: "llvm.intr.memset"(%{{.*}}, %{{.*}}, %{{.*}}, %[[falseval]]) : (!llvm.ptr, i8, i32, i1) -> ()
+  ; CHECK: "llvm.intr.memset"(%{{.*}}, %{{.*}}, %{{.*}}) <{isVolatile = false}> : (!llvm.ptr, i8, i32) -> ()
   call void @llvm.memset.p0.i32(ptr %1, i8 %2, i32 %0, i1 false)
   ret void
 }
 
 ; CHECK-LABEL:  llvm.func @sadd_with_overflow_test
 define void @sadd_with_overflow_test(i32 %0, i32 %1, <8 x i32> %2, <8 x i32> %3) {
   ; CHECK: "llvm.intr.sadd.with.overflow"(%{{.*}}, %{{.*}}) : (i32, i32) -> !llvm.struct<(i32, i1)>
   %5 = call { i32, i1 } @llvm.sadd.with.overflow.i32(i32 %0, i32 %1)
define void @vector_predication_intrinsics(<8 x i32> %0, <8 x i32> %1, <8 x float> %2, <8 x float> %3, <8 x i64> %4, <8 x double> %5, <8 x ptr> %6, i32 %7, float %8, ptr %9, ptr %10, <8 x i1> %11, i32 %12) {
   ret void
 }
 
 declare float @llvm.fmuladd.f32(float, float, float)
 declare <8 x float> @llvm.fmuladd.v8f32(<8 x float>, <8 x float>, <8 x float>)
 declare float @llvm.fma.f32(float, float, float)
 declare <8 x float> @llvm.fma.v8f32(<8 x float>, <8 x float>, <8 x float>)
 declare void @llvm.prefetch.p0(ptr nocapture readonly, i32 immarg, i32 immarg, i32)
+declare i1 @llvm.is.fpclass.f32(float, i32 immarg)
+declare <8 x i1> @llvm.is.fpclass.v8f32(<8 x float>, i32 immarg)
 declare float @llvm.exp.f32(float)
 declare <8 x float> @llvm.exp.v8f32(<8 x float>)
 declare float @llvm.exp2.f32(float)
 declare <8 x float> @llvm.exp2.v8f32(<8 x float>)
 declare float @llvm.log.f32(float)
 declare <8 x float> @llvm.log.v8f32(<8 x float>)
 declare float @llvm.log10.f32(float)
 declare <8 x float> @llvm.log10.v8f32(<8 x float>)

 declare i32 @llvm.bitreverse.i32(i32)
 declare <8 x i32> @llvm.bitreverse.v8i32(<8 x i32>)
 declare i32 @llvm.bswap.i32(i32)
 declare <8 x i32> @llvm.bswap.v8i32(<8 x i32>)
 declare i32 @llvm.ctlz.i32(i32, i1 immarg)
 declare <8 x i32> @llvm.ctlz.v8i32(<8 x i32>, i1 immarg)
 declare i32 @llvm.cttz.i32(i32, i1 immarg)
 declare <8 x i32> @llvm.cttz.v8i32(<8 x i32>, i1 immarg)
+declare i32 @llvm.abs.i32(i32, i1 immarg)
+declare <8 x i32> @llvm.abs.v8i32(<8 x i32>, i1 immarg)
 declare i32 @llvm.ctpop.i32(i32)
 declare <8 x i32> @llvm.ctpop.v8i32(<8 x i32>)
 declare i32 @llvm.fshl.i32(i32, i32, i32)
 declare <8 x i32> @llvm.fshl.v8i32(<8 x i32>, <8 x i32>, <8 x i32>)
 declare i32 @llvm.fshr.i32(i32, i32, i32)
 declare <8 x i32> @llvm.fshr.v8i32(<8 x i32>, <8 x i32>, <8 x i32>)
 declare float @llvm.maximum.f32(float, float)
 declare <8 x float> @llvm.maximum.v8f32(<8 x float>, <8 x float>)

 declare void @llvm.masked.scatter.v7f32.v7p0(<7 x float>, <7 x ptr>, i32 immarg, <7 x i1>)
 declare <7 x float> @llvm.masked.expandload.v7f32(ptr, <7 x i1>, <7 x float>)
 declare void @llvm.masked.compressstore.v7f32(<7 x float>, ptr, <7 x i1>)
 declare void @llvm.trap()
 declare void @llvm.debugtrap()
 declare void @llvm.ubsantrap(i8 immarg)
 declare void @llvm.memcpy.p0.p0.i32(ptr noalias nocapture writeonly, ptr noalias nocapture readonly, i32, i1 immarg)
 declare void @llvm.memcpy.inline.p0.p0.i64(ptr noalias nocapture writeonly, ptr noalias nocapture readonly, i64 immarg, i1 immarg)
+declare void @llvm.memcpy.inline.p0.p0.i32(ptr noalias nocapture writeonly, ptr noalias nocapture readonly, i32 immarg, i1 immarg)
 declare void @llvm.memmove.p0.p0.i32(ptr nocapture writeonly, ptr nocapture readonly, i32, i1 immarg)
 declare void @llvm.memset.p0.i32(ptr nocapture writeonly, i8, i32, i1 immarg)
 declare { i32, i1 } @llvm.sadd.with.overflow.i32(i32, i32)
 declare { <8 x i32>, <8 x i1> } @llvm.sadd.with.overflow.v8i32(<8 x i32>, <8 x i32>)
 declare { i32, i1 } @llvm.uadd.with.overflow.i32(i32, i32)
 declare { <8 x i32>, <8 x i1> } @llvm.uadd.with.overflow.v8i32(<8 x i32>, <8 x i32>)
 declare { i32, i1 } @llvm.ssub.with.overflow.i32(i32, i32)
 declare { <8 x i32>, <8 x i1> } @llvm.ssub.with.overflow.v8i32(<8 x i32>, <8 x i32>)

mlir/test/Target/LLVMIR/llvmir-intrinsics.mlir

 // RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
 
 // CHECK-LABEL: @intrinsics
 llvm.func @intrinsics(%arg0: f32, %arg1: f32, %arg2: vector<8xf32>, %arg3: !llvm.ptr<i8>) {
-  %c3 = llvm.mlir.constant(3 : i32) : i32
-  %c1 = llvm.mlir.constant(1 : i32) : i32
-  %c0 = llvm.mlir.constant(0 : i32) : i32
   // CHECK: call float @llvm.fmuladd.f32
   "llvm.intr.fmuladd"(%arg0, %arg1, %arg0) : (f32, f32, f32) -> f32
   // CHECK: call <8 x float> @llvm.fmuladd.v8f32
   "llvm.intr.fmuladd"(%arg2, %arg2, %arg2) : (vector<8xf32>, vector<8xf32>, vector<8xf32>) -> vector<8xf32>
   // CHECK: call float @llvm.fma.f32
   "llvm.intr.fma"(%arg0, %arg1, %arg0) : (f32, f32, f32) -> f32
   // CHECK: call <8 x float> @llvm.fma.v8f32
   "llvm.intr.fma"(%arg2, %arg2, %arg2) : (vector<8xf32>, vector<8xf32>, vector<8xf32>) -> vector<8xf32>
   // CHECK: call void @llvm.prefetch.p0(ptr %3, i32 0, i32 3, i32 1)
-  "llvm.intr.prefetch"(%arg3, %c0, %c3, %c1) : (!llvm.ptr<i8>, i32, i32, i32) -> ()
+  "llvm.intr.prefetch"(%arg3) <{cache = 1 : i32, hint = 3 : i32, rw = 0 : i32}> : (!llvm.ptr<i8>) -> ()
   llvm.return
 }
 
 // CHECK-LABEL: @fpclass_test
 llvm.func @fpclass_test(%arg0: f32) -> i1 {
-  %checkNan = llvm.mlir.constant(0 : i32) : i32
   // CHECK: call i1 @llvm.is.fpclass
-  %0 = "llvm.intr.is.fpclass"(%arg0, %checkNan) : (f32, i32) -> i1
+  %0 = "llvm.intr.is.fpclass"(%arg0) <{bit = 0 : i32 }>: (f32) -> i1
   llvm.return %0 : i1
 }
 
 // CHECK-LABEL: @exp_test
 llvm.func @exp_test(%arg0: f32, %arg1: vector<8xf32>) {
   // CHECK: call float @llvm.exp.f32
   "llvm.intr.exp"(%arg0) : (f32) -> f32
   // CHECK: call <8 x float> @llvm.exp.v8f32
llvm.func @byteswap_test(%arg0: i32, %arg1: vector<8xi32>) {
   "llvm.intr.bswap"(%arg0) : (i32) -> i32
   // CHECK: call <8 x i32> @llvm.bswap.v8i32
   "llvm.intr.bswap"(%arg1) : (vector<8xi32>) -> vector<8xi32>
   llvm.return
 }
 
 // CHECK-LABEL: @ctlz_test
 llvm.func @ctlz_test(%arg0: i32, %arg1: vector<8xi32>) {
-  %i1 = llvm.mlir.constant(false) : i1
   // CHECK: call i32 @llvm.ctlz.i32
-  "llvm.intr.ctlz"(%arg0, %i1) : (i32, i1) -> i32
+  "llvm.intr.ctlz"(%arg0) <{is_zero_poison = 0 : i1}> : (i32) -> i32
   // CHECK: call <8 x i32> @llvm.ctlz.v8i32
-  "llvm.intr.ctlz"(%arg1, %i1) : (vector<8xi32>, i1) -> vector<8xi32>
+  "llvm.intr.ctlz"(%arg1) <{is_zero_poison = 1 : i1}> : (vector<8xi32>) -> vector<8xi32>
   llvm.return
 }
 
 // CHECK-LABEL: @cttz_test
 llvm.func @cttz_test(%arg0: i32, %arg1: vector<8xi32>) {
-  %i1 = llvm.mlir.constant(false) : i1
   // CHECK: call i32 @llvm.cttz.i32
-  "llvm.intr.cttz"(%arg0, %i1) : (i32, i1) -> i32
+  "llvm.intr.cttz"(%arg0) <{is_zero_poison = 0 : i1}> : (i32) -> i32
   // CHECK: call <8 x i32> @llvm.cttz.v8i32
-  "llvm.intr.cttz"(%arg1, %i1) : (vector<8xi32>, i1) -> vector<8xi32>
+  "llvm.intr.cttz"(%arg1) <{is_zero_poison = 1 : i1}> : (vector<8xi32>) -> vector<8xi32>
+  llvm.return
+}
+
+// CHECK-LABEL: @abs_test
+llvm.func @abs_test(%arg0: i32, %arg1: vector<8xi32>) {
+  // CHECK: call i32 @llvm.abs.i32
+  "llvm.intr.abs"(%arg0) <{is_int_min_poison = 0 : i1}> : (i32) -> i32
+  // CHECK: call <8 x i32> @llvm.abs.v8i32
+  "llvm.intr.abs"(%arg1) <{is_int_min_poison = 1 : i1}> : (vector<8xi32>) -> vector<8xi32>
   llvm.return
 }
 
 // CHECK-LABEL: @ctpop_test
 llvm.func @ctpop_test(%arg0: i32, %arg1: vector<8xi32>) {
   // CHECK: call i32 @llvm.ctpop.i32
   "llvm.intr.ctpop"(%arg0) : (i32) -> i32
   // CHECK: call <8 x i32> @llvm.ctpop.v8i32
llvm.func @trap_intrinsics() {
   "llvm.intr.debugtrap"() : () -> ()
   // CHECK: call void @llvm.ubsantrap(i8 1)
   "llvm.intr.ubsantrap"() {failureKind = 1 : i8} : () -> ()
   llvm.return
 }
 
 // CHECK-LABEL: @memcpy_test
 llvm.func @memcpy_test(%arg0: i32, %arg2: !llvm.ptr<i8>, %arg3: !llvm.ptr<i8>) {
-  %i1 = llvm.mlir.constant(false) : i1
-  // CHECK: call void @llvm.memcpy.p0.p0.i32(ptr %{{.*}}, ptr %{{.*}}, i32 %{{.*}}, i1 {{.*}})
-  "llvm.intr.memcpy"(%arg2, %arg3, %arg0, %i1) : (!llvm.ptr<i8>, !llvm.ptr<i8>, i32, i1) -> ()
-  %sz = llvm.mlir.constant(10: i64) : i64
-  // CHECK: call void @llvm.memcpy.inline.p0.p0.i64(ptr %{{.*}}, ptr %{{.*}}, i64 10, i1 {{.*}})
-  "llvm.intr.memcpy.inline"(%arg2, %arg3, %sz, %i1) : (!llvm.ptr<i8>, !llvm.ptr<i8>, i64, i1) -> ()
+  // CHECK: call void @llvm.memcpy.p0.p0.i32(ptr %{{.*}}, ptr %{{.*}}, i32 %{{.*}}, i1 false
+  "llvm.intr.memcpy"(%arg2, %arg3, %arg0) <{isVolatile = false}> : (!llvm.ptr<i8>, !llvm.ptr<i8>, i32) -> ()
+  // CHECK: call void @llvm.memcpy.inline.p0.p0.i32(ptr %{{.*}}, ptr %{{.*}}, i32 10, i1 true
+  "llvm.intr.memcpy.inline"(%arg2, %arg3) <{isVolatile = true, len = 10 : i32}> : (!llvm.ptr<i8>, !llvm.ptr<i8>) -> ()
+  // CHECK: call void @llvm.memcpy.inline.p0.p0.i64(ptr %{{.*}}, ptr %{{.*}}, i64 10, i1 true
+  "llvm.intr.memcpy.inline"(%arg2, %arg3) <{isVolatile = true, len = 10 : i64}> : (!llvm.ptr<i8>, !llvm.ptr<i8>) -> ()
   llvm.return
 }
 
 // CHECK-LABEL: @memmove_test
 llvm.func @memmove_test(%arg0: i32, %arg2: !llvm.ptr<i8>, %arg3: !llvm.ptr<i8>) {
-  %i1 = llvm.mlir.constant(false) : i1
-  // CHECK: call void @llvm.memmove.p0.p0.i32(ptr %{{.*}}, ptr %{{.*}}, i32 %{{.*}}, i1 {{.*}})
-  "llvm.intr.memmove"(%arg2, %arg3, %arg0, %i1) : (!llvm.ptr<i8>, !llvm.ptr<i8>, i32, i1) -> ()
+  // CHECK: call void @llvm.memmove.p0.p0.i32(ptr %{{.*}}, ptr %{{.*}}, i32 %{{.*}}, i1 false
+  "llvm.intr.memmove"(%arg2, %arg3, %arg0) <{isVolatile = false}> : (!llvm.ptr<i8>, !llvm.ptr<i8>, i32) -> ()
   llvm.return
 }
 
 // CHECK-LABEL: @memset_test
 llvm.func @memset_test(%arg0: i32, %arg2: !llvm.ptr<i8>, %arg3: i8) {
   %i1 = llvm.mlir.constant(false) : i1
-  // CHECK: call void @llvm.memset.p0.i32(ptr %{{.*}}, i8 %{{.*}}, i32 %{{.*}}, i1 {{.*}})
-  "llvm.intr.memset"(%arg2, %arg3, %arg0, %i1) : (!llvm.ptr<i8>, i8, i32, i1) -> ()
+  // CHECK: call void @llvm.memset.p0.i32(ptr %{{.*}}, i8 %{{.*}}, i32 %{{.*}}, i1 false
+  "llvm.intr.memset"(%arg2, %arg3, %arg0) <{isVolatile = false}> : (!llvm.ptr<i8>, i8, i32) -> ()
   llvm.return
 }
 
 // CHECK-LABEL: @sadd_with_overflow_test
 llvm.func @sadd_with_overflow_test(%arg0: i32, %arg1: i32, %arg2: vector<8xi32>, %arg3: vector<8xi32>) {
   // CHECK: call { i32, i1 } @llvm.sadd.with.overflow.i32
   "llvm.intr.sadd.with.overflow"(%arg0, %arg1) : (i32, i32) -> !llvm.struct<(i32, i1)>
   // CHECK: call { <8 x i32>, <8 x i1> } @llvm.sadd.with.overflow.v8i32

 }
 
 // Check that intrinsics are declared with appropriate types.
 // CHECK-DAG: declare float @llvm.fma.f32(float, float, float)
 // CHECK-DAG: declare <8 x float> @llvm.fma.v8f32(<8 x float>, <8 x float>, <8 x float>) #0
 // CHECK-DAG: declare float @llvm.fmuladd.f32(float, float, float)
 // CHECK-DAG: declare <8 x float> @llvm.fmuladd.v8f32(<8 x float>, <8 x float>, <8 x float>) #0
 // CHECK-DAG: declare void @llvm.prefetch.p0(ptr nocapture readonly, i32 immarg, i32 immarg, i32 immarg)
+// CHECK-DAG: declare i1 @llvm.is.fpclass.f32(float, i32 immarg)
 // CHECK-DAG: declare float @llvm.exp.f32(float)
 // CHECK-DAG: declare <8 x float> @llvm.exp.v8f32(<8 x float>) #0
 // CHECK-DAG: declare float @llvm.log.f32(float)
 // CHECK-DAG: declare <8 x float> @llvm.log.v8f32(<8 x float>) #0
 // CHECK-DAG: declare float @llvm.log10.f32(float)
 // CHECK-DAG: declare <8 x float> @llvm.log10.v8f32(<8 x float>) #0
 // CHECK-DAG: declare float @llvm.log2.f32(float)
 // CHECK-DAG: declare <8 x float> @llvm.log2.v8f32(<8 x float>) #0

 // CHECK-DAG: declare <7 x float> @llvm.masked.gather.v7f32.v7p0(<7 x ptr>, i32 immarg, <7 x i1>, <7 x float>)
 // CHECK-DAG: declare void @llvm.masked.scatter.v7f32.v7p0(<7 x float>, <7 x ptr>, i32 immarg, <7 x i1>)
 // CHECK-DAG: declare <7 x float> @llvm.masked.expandload.v7f32(ptr nocapture, <7 x i1>, <7 x float>)
 // CHECK-DAG: declare void @llvm.masked.compressstore.v7f32(<7 x float>, ptr nocapture, <7 x i1>)
 // CHECK-DAG: declare void @llvm.trap()
 // CHECK-DAG: declare void @llvm.debugtrap()
 // CHECK-DAG: declare void @llvm.ubsantrap(i8 immarg)
 // CHECK-DAG: declare void @llvm.memcpy.p0.p0.i32(ptr noalias nocapture writeonly, ptr noalias nocapture readonly, i32, i1 immarg)
+// CHECK-DAG: declare void @llvm.memcpy.inline.p0.p0.i32(ptr noalias nocapture writeonly, ptr noalias nocapture readonly, i32 immarg, i1 immarg)
 // CHECK-DAG: declare void @llvm.memcpy.inline.p0.p0.i64(ptr noalias nocapture writeonly, ptr noalias nocapture readonly, i64 immarg, i1 immarg)
 // CHECK-DAG: declare { i32, i1 } @llvm.sadd.with.overflow.i32(i32, i32)
 // CHECK-DAG: declare { <8 x i32>, <8 x i1> } @llvm.sadd.with.overflow.v8i32(<8 x i32>, <8 x i32>)
 // CHECK-DAG: declare { i32, i1 } @llvm.uadd.with.overflow.i32(i32, i32)
 // CHECK-DAG: declare { <8 x i32>, <8 x i1> } @llvm.uadd.with.overflow.v8i32(<8 x i32>, <8 x i32>)
 // CHECK-DAG: declare { i32, i1 } @llvm.ssub.with.overflow.i32(i32, i32)
 // CHECK-DAG: declare { <8 x i32>, <8 x i1> } @llvm.ssub.with.overflow.v8i32(<8 x i32>, <8 x i32>)
 // CHECK-DAG: declare { i32, i1 } @llvm.usub.with.overflow.i32(i32, i32)

mlir/test/Target/LLVMIR/llvmir-invalid.mlir

 llvm.func @powi_intr_wrong_type(%arg0 : f32, %arg1 : f32) -> f32 {
   // expected-error @below{{op operand #1 must be signless integer, but got 'f32'}}
   %0 = "llvm.intr.powi"(%arg0, %arg1) : (f32, f32) -> f32
   llvm.return %0 : f32
 }
 
 // -----
 
-llvm.func @ctlz_intr_wrong_type(%arg0 : i32, %arg1 : i32) -> i32 {
-  // expected-error @below{{op operand #1 must be 1-bit signless integer, but got 'i32'}}
-  %0 = "llvm.intr.ctlz"(%arg0, %arg1) : (i32, i32) -> i32
-  llvm.return %0 : i32
-}
-
-// -----
-
-llvm.func @memcpy_intr_wrong_type(%src : i64, %dst : i64, %len : i64, %volatile : i1) {
+llvm.func @memcpy_intr_wrong_type(%src : i64, %dst : i64, %len : i64) {
   // expected-error @below{{op operand #0 must be LLVM pointer type, but got 'i64'}}
-  "llvm.intr.memcpy"(%src, %dst, %len, %volatile) : (i64, i64, i64, i1) -> ()
-  llvm.return
-}
-
-// -----
-
-llvm.func @memcpy_inline_intr_wrong_type(%src : !llvm.ptr, %dst : !llvm.ptr, %len : i64, %volatile : i32) {
-  // expected-error @below{{op operand #3 must be 1-bit signless integer, but got 'i32'}}
-  "llvm.intr.memcpy.inline"(%src, %dst, %len, %volatile) : (!llvm.ptr, !llvm.ptr, i64, i32) -> ()
+  "llvm.intr.memcpy"(%src, %dst, %len) <{isVolatile = false}> : (i64, i64, i64) -> ()
   llvm.return
 }
 
 // -----
 
-llvm.func @memmove_intr_wrong_type(%src : !llvm.ptr, %dst : i64, %len : i64, %volatile : i1) {
+llvm.func @memmove_intr_wrong_type(%src : !llvm.ptr, %dst : i64, %len : i64) {
   // expected-error @below{{op operand #1 must be LLVM pointer type, but got 'i64'}}
-  "llvm.intr.memmove"(%src, %dst, %len, %volatile) : (!llvm.ptr, i64, i64, i1) -> ()
+  "llvm.intr.memmove"(%src, %dst, %len) <{isVolatile = false}> : (!llvm.ptr, i64, i64) -> ()
   llvm.return
 }
 
 // -----
 
-llvm.func @memset_intr_wrong_type(%dst : !llvm.ptr, %val : i32, %len : i64, %volatile : i1) {
+llvm.func @memset_intr_wrong_type(%dst : !llvm.ptr, %val : i32, %len : i64) {
   // expected-error @below{{op operand #1 must be 8-bit signless integer, but got 'i32'}}
-  "llvm.intr.memset"(%dst, %val, %len, %volatile) : (!llvm.ptr, i32, i64, i1) -> ()
+  "llvm.intr.memset"(%dst, %val, %len) <{isVolatile = false}> : (!llvm.ptr, i32, i64) -> ()
   llvm.return
 }
 
 // -----
 
 llvm.func @sadd_overflow_intr_wrong_type(%arg0 : i32, %arg1 : f32) -> !llvm.struct<(i32, i1)> {
   // expected-error @below{{op operand #1 must be signless integer or LLVM dialect-compatible vector of signless integer, but got 'f32'}}
   %0 = "llvm.intr.sadd.with.overflow"(%arg0, %arg1) : (i32, f32) -> !llvm.struct<(i32, i1)>

mlir/test/Target/LLVMIR/llvmir.mlir

llvm.func @aliasScope(%arg1 : !llvm.ptr) {
   // CHECK:  store {{.*}}, !alias.scope ![[SCOPES1]], !noalias ![[SCOPES23:[0-9]+]]
   llvm.store %0, %arg1 {alias_scopes = [@metadata::@scope1], noalias_scopes = [@metadata::@scope2, @metadata::@scope3]} : i32, !llvm.ptr
   // CHECK:  load {{.*}}, !alias.scope ![[SCOPES2:[0-9]+]], !noalias ![[SCOPES13:[0-9]+]]
   %1 = llvm.load %arg1 {alias_scopes = [@metadata::@scope2], noalias_scopes = [@metadata::@scope1, @metadata::@scope3]} : !llvm.ptr -> i32
   // CHECK:  atomicrmw {{.*}}, !alias.scope ![[SCOPES3:[0-9]+]], !noalias ![[SCOPES12:[0-9]+]]
   %2 = llvm.atomicrmw add %arg1, %0 monotonic {alias_scopes = [@metadata::@scope3], noalias_scopes = [@metadata::@scope1, @metadata::@scope2]} : !llvm.ptr, i32
   // CHECK:  cmpxchg {{.*}}, !alias.scope ![[SCOPES3]]
   %3 = llvm.cmpxchg %arg1, %1, %2 acq_rel monotonic {alias_scopes = [@metadata::@scope3]} : !llvm.ptr, i32
-  %4 = llvm.mlir.constant(0 : i1) : i1
   %5 = llvm.mlir.constant(42 : i8) : i8
   // CHECK:  llvm.memcpy{{.*}}, !alias.scope ![[SCOPES3]]
-  "llvm.intr.memcpy"(%arg1, %arg1, %0, %4) {alias_scopes = [@metadata::@scope3]} : (!llvm.ptr, !llvm.ptr, i32, i1) -> ()
+  "llvm.intr.memcpy"(%arg1, %arg1, %0) <{isVolatile = false}> {alias_scopes = [@metadata::@scope3]} : (!llvm.ptr, !llvm.ptr, i32) -> ()
   // CHECK:  llvm.memset{{.*}}, !noalias ![[SCOPES3]]
-  "llvm.intr.memset"(%arg1, %5, %0, %4) {noalias_scopes = [@metadata::@scope3]} : (!llvm.ptr, i8, i32, i1) -> ()
+  "llvm.intr.memset"(%arg1, %5, %0) <{isVolatile = false}> {noalias_scopes = [@metadata::@scope3]} : (!llvm.ptr, i8, i32) -> ()
   // CHECK: call void @foo({{.*}} !alias.scope ![[SCOPES3]]
   llvm.call @foo(%arg1) {alias_scopes = [@metadata::@scope3]} : (!llvm.ptr) -> ()
   // CHECK: call void @foo({{.*}} !noalias ![[SCOPES3]]
   llvm.call @foo(%arg1) {noalias_scopes = [@metadata::@scope3]} : (!llvm.ptr) -> ()
   llvm.return
 }
 
 llvm.metadata @metadata {

mlir/test/Target/LLVMIR/loop-metadata.mlir

^bb4:
     // CHECK: = load i32, ptr %{{.*}} !llvm.access.group ![[ACCESS_GROUPS_NODE:[0-9]+]]
     %5 = llvm.load %4 {access_groups = [@metadata::@group1, @metadata::@group2]} : !llvm.ptr -> i32
     // CHECK: store i32 %{{.*}}, ptr %{{.*}} !llvm.access.group ![[ACCESS_GROUPS_NODE]]
     llvm.store %5, %4 {access_groups = [@metadata::@group1, @metadata::@group2]} : i32, !llvm.ptr
     // CHECK: = atomicrmw add ptr %{{.*}}, i32 %{{.*}} !llvm.access.group ![[ACCESS_GROUPS_NODE]]
     %6 = llvm.atomicrmw add %4, %5 monotonic {access_groups = [@metadata::@group1, @metadata::@group2]} : !llvm.ptr, i32
     // CHECK: = cmpxchg ptr %{{.*}}, i32 %{{.*}}, i32 %{{.*}} !llvm.access.group ![[ACCESS_GROUPS_NODE]]
     %7 = llvm.cmpxchg %4, %5, %6 acq_rel monotonic {access_groups = [@metadata::@group1, @metadata::@group2]} : !llvm.ptr, i32
-    %8 = llvm.mlir.constant(0 : i1) : i1
     %9 = llvm.mlir.constant(42 : i8) : i8
     // CHECK: llvm.memcpy{{.*}} !llvm.access.group ![[ACCESS_GROUPS_NODE]]
-    "llvm.intr.memcpy"(%4, %4, %0, %8) {access_groups = [@metadata::@group1, @metadata::@group2]} : (!llvm.ptr, !llvm.ptr, i32, i1) -> ()
+    "llvm.intr.memcpy"(%4, %4, %0) <{isVolatile = false}> {access_groups = [@metadata::@group1, @metadata::@group2]} : (!llvm.ptr, !llvm.ptr, i32) -> ()
     // CHECK: llvm.memset{{.*}} !llvm.access.group ![[ACCESS_GROUPS_NODE]]
-    "llvm.intr.memset"(%4, %9, %0, %8) {access_groups = [@metadata::@group1, @metadata::@group2]} : (!llvm.ptr, i8, i32, i1) -> ()
+    "llvm.intr.memset"(%4, %9, %0) <{isVolatile = false}> {access_groups = [@metadata::@group1, @metadata::@group2]} : (!llvm.ptr, i8, i32) -> ()
     // CHECK: call void @foo({{.*}} !llvm.access.group ![[ACCESS_GROUPS_NODE]]
     llvm.call @foo(%arg1) {access_groups = [@metadata::@group1, @metadata::@group2]} : (i32) -> ()
     // CHECK: br label {{.*}} !llvm.loop ![[LOOP_NODE]]
     llvm.br ^bb3(%3 : i32) {loop_annotation = #llvm.loop_annotation<
           licm = <disable = true>,
           interleave = <count = 1>,
           unroll = <disable = true>, pipeline = <disable = true, initiationinterval = 2>,
           parallelAccesses = @metadata::@group1, @metadata::@group2>}

mlir/test/Target/LLVMIR/tbaa.mlir

llvm.func @tbaa2(%arg0: !llvm.ptr, %arg1: !llvm.ptr) {
     %4 = llvm.fptosi %3 : f32 to i32
     %5 = llvm.getelementptr inbounds %arg0[%0, 0] : (!llvm.ptr, i32) -> !llvm.ptr, !llvm.struct<"struct.agg1_t", (i32, i32)>
     // CHECK: store i32 %{{.*}}, ptr %{{.*}},{{.*}}!tbaa ![[STAG:[0-9]*]]
     llvm.store %4, %5 {tbaa = [@__tbaa::@tbaa_tag_7]} : i32, !llvm.ptr
     // CHECK: atomicrmw add ptr %{{.*}}, i32 %{{.*}} !tbaa ![[STAG]]
     %6 = llvm.atomicrmw add %5, %4 monotonic {tbaa = [@__tbaa::@tbaa_tag_7]} : !llvm.ptr, i32
     // CHECK: cmpxchg ptr %{{.*}}, i32 %{{.*}}, i32 %{{.*}} !tbaa ![[STAG]]
     %7 = llvm.cmpxchg %5, %6, %4 acq_rel monotonic {tbaa = [@__tbaa::@tbaa_tag_7]} : !llvm.ptr, i32
-    %8 = llvm.mlir.constant(0 : i1) : i1
     %9 = llvm.mlir.constant(42 : i8) : i8
     // CHECK: llvm.memcpy{{.*}} !tbaa ![[STAG]]
-    "llvm.intr.memcpy"(%arg1, %arg1, %0, %8) {tbaa = [@__tbaa::@tbaa_tag_7]} : (!llvm.ptr, !llvm.ptr, i32, i1) -> ()
+    "llvm.intr.memcpy"(%arg1, %arg1, %0) <{isVolatile = false}> {tbaa = [@__tbaa::@tbaa_tag_7]} : (!llvm.ptr, !llvm.ptr, i32) -> ()
     // CHECK: llvm.memset{{.*}} !tbaa ![[STAG]]
-    "llvm.intr.memset"(%arg1, %9, %0, %8) {tbaa = [@__tbaa::@tbaa_tag_7]} : (!llvm.ptr, i8, i32, i1) -> ()
+    "llvm.intr.memset"(%arg1, %9, %0) <{isVolatile = false}> {tbaa = [@__tbaa::@tbaa_tag_7]} : (!llvm.ptr, i8, i32) -> ()
     // CHECK: call void @foo({{.*}} !tbaa ![[STAG]]
     llvm.call @foo(%arg1) {tbaa = [@__tbaa::@tbaa_tag_7]} : (!llvm.ptr) -> ()
     llvm.return
   }
 }
 
 // CHECK-DAG: ![[LTAG]] = !{![[AGG2T:[0-9]*]], ![[F32T:[0-9]*]], i64 0}
 // CHECK-DAG: ![[AGG2T]] = !{!"agg2_t", ![[F32T]], i64 0, ![[F32T]], i64 4}
 // CHECK-DAG: ![[I64T]] = !{!"float", ![[CHART:[0-9]*]], i64 0}
 // CHECK-DAG: ![[CHART]] = !{!"omnipotent char", ![[ROOT:[0-9]*]], i64 0}
 // CHECK-DAG: ![[ROOT]] = !{!"Simple C/C++ TBAA"}
 // CHECK-DAG: ![[STAG]] = !{![[AGG1T:[0-9]*]], ![[I32T:[0-9]*]], i64 0}
 // CHECK-DAG: ![[AGG1T]] = !{!"agg1_t", ![[I32T]], i64 0, ![[I32T]], i64 4}
 // CHECK-DAG: ![[I32T]] = !{!"int", ![[CHART]], i64 0}