NFC: Replace ValuePtr with Value and remove it now that Value is value-typed.

ValuePtr was a temporary typedef during the transition to a value-typed Value.

PiperOrigin-RevId: 286945714

1 files changed, 24 insertions, 25 deletions

diff --git a/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp b/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp
index 3383cf13d36..19dabcdafee 100644
--- a/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp
+++ b/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp
@@ -105,7 +105,7 @@ private:
   }
 
   // Allocate a void pointer on the stack.
-  ValuePtr allocatePointer(OpBuilder &builder, Location loc) {
+  Value allocatePointer(OpBuilder &builder, Location loc) {
     auto one = builder.create<LLVM::ConstantOp>(loc, getInt32Type(),
                                                 builder.getI32IntegerAttr(1));
     return builder.create<LLVM::AllocaOp>(loc, getPointerPointerType(), one,
@@ -113,9 +113,9 @@ private:
   }
 
   void declareCudaFunctions(Location loc);
-  ValuePtr setupParamsArray(gpu::LaunchFuncOp launchOp, OpBuilder &builder);
-  ValuePtr generateKernelNameConstant(StringRef name, Location loc,
-                                      OpBuilder &builder);
+  Value setupParamsArray(gpu::LaunchFuncOp launchOp, OpBuilder &builder);
+  Value generateKernelNameConstant(StringRef name, Location loc,
+                                   OpBuilder &builder);
   void translateGpuLaunchCalls(mlir::gpu::LaunchFuncOp launchOp);
 
 public:
@@ -239,9 +239,8 @@ void GpuLaunchFuncToCudaCallsPass::declareCudaFunctions(Location loc) {
 // for (i : [0, NumKernelOperands))
 //   %array[i] = cast<void*>(KernelOperand[i])
 // return %array
-ValuePtr
-GpuLaunchFuncToCudaCallsPass::setupParamsArray(gpu::LaunchFuncOp launchOp,
-                                               OpBuilder &builder) {
+Value GpuLaunchFuncToCudaCallsPass::setupParamsArray(gpu::LaunchFuncOp launchOp,
+                                                     OpBuilder &builder) {
   auto numKernelOperands = launchOp.getNumKernelOperands();
   Location loc = launchOp.getLoc();
   auto one = builder.create<LLVM::ConstantOp>(loc, getInt32Type(),
@@ -255,7 +254,7 @@ GpuLaunchFuncToCudaCallsPass::setupParamsArray(gpu::LaunchFuncOp launchOp,
   for (unsigned idx = 0; idx < numKernelOperands; ++idx) {
     auto operand = launchOp.getKernelOperand(idx);
     auto llvmType = operand->getType().cast<LLVM::LLVMType>();
-    ValuePtr memLocation = builder.create<LLVM::AllocaOp>(
+    Value memLocation = builder.create<LLVM::AllocaOp>(
         loc, llvmType.getPointerTo(), one, /*alignment=*/1);
     builder.create<LLVM::StoreOp>(loc, operand, memLocation);
     auto casted =
@@ -271,12 +270,12 @@ GpuLaunchFuncToCudaCallsPass::setupParamsArray(gpu::LaunchFuncOp launchOp,
           getModule().lookupSymbol<LLVM::LLVMFuncOp>(kMcuMemHostRegister);
       auto nullPtr = builder.create<LLVM::NullOp>(loc, llvmType.getPointerTo());
       auto gep = builder.create<LLVM::GEPOp>(loc, llvmType.getPointerTo(),
-                                             ArrayRef<ValuePtr>{nullPtr, one});
+                                             ArrayRef<Value>{nullPtr, one});
       auto size = builder.create<LLVM::PtrToIntOp>(loc, getInt64Type(), gep);
       builder.create<LLVM::CallOp>(loc, ArrayRef<Type>{},
                                    builder.getSymbolRefAttr(registerFunc),
-                                   ArrayRef<ValuePtr>{casted, size});
-      ValuePtr memLocation = builder.create<LLVM::AllocaOp>(
+                                   ArrayRef<Value>{casted, size});
+      Value memLocation = builder.create<LLVM::AllocaOp>(
           loc, getPointerPointerType(), one, /*alignment=*/1);
       builder.create<LLVM::StoreOp>(loc, casted, memLocation);
       casted =
@@ -286,7 +285,7 @@ GpuLaunchFuncToCudaCallsPass::setupParamsArray(gpu::LaunchFuncOp launchOp,
     auto index = builder.create<LLVM::ConstantOp>(
         loc, getInt32Type(), builder.getI32IntegerAttr(idx));
     auto gep = builder.create<LLVM::GEPOp>(loc, getPointerPointerType(), array,
-                                           ArrayRef<ValuePtr>{index});
+                                           ArrayRef<Value>{index});
     builder.create<LLVM::StoreOp>(loc, casted, gep);
   }
   return array;
@@ -302,7 +301,7 @@ GpuLaunchFuncToCudaCallsPass::setupParamsArray(gpu::LaunchFuncOp launchOp,
 //   %1 = llvm.constant (0 : index)
 //   %2 = llvm.getelementptr %0[%1, %1] : !llvm<"i8*">
 // }
-ValuePtr GpuLaunchFuncToCudaCallsPass::generateKernelNameConstant(
+Value GpuLaunchFuncToCudaCallsPass::generateKernelNameConstant(
     StringRef name, Location loc, OpBuilder &builder) {
   // Make sure the trailing zero is included in the constant.
   std::vector<char> kernelName(name.begin(), name.end());
@@ -358,7 +357,7 @@ void GpuLaunchFuncToCudaCallsPass::translateGpuLaunchCalls(
   assert(kernelModule.getName() && "expected a named module");
   SmallString<128> nameBuffer(*kernelModule.getName());
   nameBuffer.append(kCubinStorageSuffix);
-  ValuePtr data = LLVM::createGlobalString(
+  Value data = LLVM::createGlobalString(
       loc, builder, nameBuffer.str(), cubinAttr.getValue(),
       LLVM::Linkage::Internal, getLLVMDialect());
 
@@ -369,7 +368,7 @@ void GpuLaunchFuncToCudaCallsPass::translateGpuLaunchCalls(
       getModule().lookupSymbol<LLVM::LLVMFuncOp>(cuModuleLoadName);
   builder.create<LLVM::CallOp>(loc, ArrayRef<Type>{getCUResultType()},
                                builder.getSymbolRefAttr(cuModuleLoad),
-                               ArrayRef<ValuePtr>{cuModule, data});
+                               ArrayRef<Value>{cuModule, data});
   // Get the function from the module. The name corresponds to the name of
   // the kernel function.
   auto cuOwningModuleRef =
@@ -381,13 +380,13 @@ void GpuLaunchFuncToCudaCallsPass::translateGpuLaunchCalls(
   builder.create<LLVM::CallOp>(
       loc, ArrayRef<Type>{getCUResultType()},
       builder.getSymbolRefAttr(cuModuleGetFunction),
-      ArrayRef<ValuePtr>{cuFunction, cuOwningModuleRef, kernelName});
+      ArrayRef<Value>{cuFunction, cuOwningModuleRef, kernelName});
   // Grab the global stream needed for execution.
   auto cuGetStreamHelper =
       getModule().lookupSymbol<LLVM::LLVMFuncOp>(cuGetStreamHelperName);
   auto cuStream = builder.create<LLVM::CallOp>(
       loc, ArrayRef<Type>{getPointerType()},
-      builder.getSymbolRefAttr(cuGetStreamHelper), ArrayRef<ValuePtr>{});
+      builder.getSymbolRefAttr(cuGetStreamHelper), ArrayRef<Value>{});
   // Invoke the function with required arguments.
   auto cuLaunchKernel =
       getModule().lookupSymbol<LLVM::LLVMFuncOp>(cuLaunchKernelName);
@@ -399,19 +398,19 @@ void GpuLaunchFuncToCudaCallsPass::translateGpuLaunchCalls(
   builder.create<LLVM::CallOp>(
       loc, ArrayRef<Type>{getCUResultType()},
       builder.getSymbolRefAttr(cuLaunchKernel),
-      ArrayRef<ValuePtr>{cuFunctionRef, launchOp.getOperand(0),
-                         launchOp.getOperand(1), launchOp.getOperand(2),
-                         launchOp.getOperand(3), launchOp.getOperand(4),
-                         launchOp.getOperand(5), zero, /* sharedMemBytes */
-                         cuStream.getResult(0),        /* stream */
-                         paramsArray,                  /* kernel params */
-                         nullpointer /* extra */});
+      ArrayRef<Value>{cuFunctionRef, launchOp.getOperand(0),
+                      launchOp.getOperand(1), launchOp.getOperand(2),
+                      launchOp.getOperand(3), launchOp.getOperand(4),
+                      launchOp.getOperand(5), zero, /* sharedMemBytes */
+                      cuStream.getResult(0),        /* stream */
+                      paramsArray,                  /* kernel params */
+                      nullpointer /* extra */});
   // Sync on the stream to make it synchronous.
   auto cuStreamSync =
       getModule().lookupSymbol<LLVM::LLVMFuncOp>(cuStreamSynchronizeName);
   builder.create<LLVM::CallOp>(loc, ArrayRef<Type>{getCUResultType()},
                                builder.getSymbolRefAttr(cuStreamSync),
-                               ArrayRef<ValuePtr>(cuStream.getResult(0)));
+                               ArrayRef<Value>(cuStream.getResult(0)));
   launchOp.erase();
 }