diff options
Diffstat (limited to 'mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp')
| -rw-r--r-- | mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp | 46 |
1 files changed, 23 insertions, 23 deletions
diff --git a/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp b/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp index f342083bee7..840ad6ba701 100644 --- a/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp +++ b/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp @@ -114,7 +114,7 @@ private: } // Allocate a void pointer on the stack. - Value *allocatePointer(OpBuilder &builder, Location loc) { + ValuePtr allocatePointer(OpBuilder &builder, Location loc) { auto one = builder.create<LLVM::ConstantOp>(loc, getInt32Type(), builder.getI32IntegerAttr(1)); return builder.create<LLVM::AllocaOp>(loc, getPointerPointerType(), one, @@ -122,9 +122,9 @@ private: } void declareCudaFunctions(Location loc); - Value *setupParamsArray(gpu::LaunchFuncOp launchOp, OpBuilder &builder); - Value *generateKernelNameConstant(StringRef name, Location loc, - OpBuilder &builder); + ValuePtr setupParamsArray(gpu::LaunchFuncOp launchOp, OpBuilder &builder); + ValuePtr generateKernelNameConstant(StringRef name, Location loc, + OpBuilder &builder); void translateGpuLaunchCalls(mlir::gpu::LaunchFuncOp launchOp); public: @@ -248,7 +248,7 @@ void GpuLaunchFuncToCudaCallsPass::declareCudaFunctions(Location loc) { // for (i : [0, NumKernelOperands)) // %array[i] = cast<void*>(KernelOperand[i]) // return %array -Value * +ValuePtr GpuLaunchFuncToCudaCallsPass::setupParamsArray(gpu::LaunchFuncOp launchOp, OpBuilder &builder) { auto numKernelOperands = launchOp.getNumKernelOperands(); @@ -264,7 +264,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>(); - Value *memLocation = builder.create<LLVM::AllocaOp>( + ValuePtr memLocation = builder.create<LLVM::AllocaOp>( loc, llvmType.getPointerTo(), one, /*alignment=*/1); builder.create<LLVM::StoreOp>(loc, operand, memLocation); auto casted = @@ -280,12 +280,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<Value *>{nullPtr, one}); + ArrayRef<ValuePtr>{nullPtr, one}); auto size = builder.create<LLVM::PtrToIntOp>(loc, getInt64Type(), gep); builder.create<LLVM::CallOp>(loc, ArrayRef<Type>{}, builder.getSymbolRefAttr(registerFunc), - ArrayRef<Value *>{casted, size}); - Value *memLocation = builder.create<LLVM::AllocaOp>( + ArrayRef<ValuePtr>{casted, size}); + ValuePtr memLocation = builder.create<LLVM::AllocaOp>( loc, getPointerPointerType(), one, /*alignment=*/1); builder.create<LLVM::StoreOp>(loc, casted, memLocation); casted = @@ -295,7 +295,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<Value *>{index}); + ArrayRef<ValuePtr>{index}); builder.create<LLVM::StoreOp>(loc, casted, gep); } return array; @@ -311,7 +311,7 @@ GpuLaunchFuncToCudaCallsPass::setupParamsArray(gpu::LaunchFuncOp launchOp, // %1 = llvm.constant (0 : index) // %2 = llvm.getelementptr %0[%1, %1] : !llvm<"i8*"> // } -Value *GpuLaunchFuncToCudaCallsPass::generateKernelNameConstant( +ValuePtr 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()); @@ -367,7 +367,7 @@ void GpuLaunchFuncToCudaCallsPass::translateGpuLaunchCalls( assert(kernelModule.getName() && "expected a named module"); SmallString<128> nameBuffer(*kernelModule.getName()); nameBuffer.append(kCubinStorageSuffix); - Value *data = LLVM::createGlobalString( + ValuePtr data = LLVM::createGlobalString( loc, builder, nameBuffer.str(), cubinAttr.getValue(), LLVM::Linkage::Internal, getLLVMDialect()); @@ -378,7 +378,7 @@ void GpuLaunchFuncToCudaCallsPass::translateGpuLaunchCalls( getModule().lookupSymbol<LLVM::LLVMFuncOp>(cuModuleLoadName); builder.create<LLVM::CallOp>(loc, ArrayRef<Type>{getCUResultType()}, builder.getSymbolRefAttr(cuModuleLoad), - ArrayRef<Value *>{cuModule, data}); + ArrayRef<ValuePtr>{cuModule, data}); // Get the function from the module. The name corresponds to the name of // the kernel function. auto cuOwningModuleRef = @@ -390,13 +390,13 @@ void GpuLaunchFuncToCudaCallsPass::translateGpuLaunchCalls( builder.create<LLVM::CallOp>( loc, ArrayRef<Type>{getCUResultType()}, builder.getSymbolRefAttr(cuModuleGetFunction), - ArrayRef<Value *>{cuFunction, cuOwningModuleRef, kernelName}); + ArrayRef<ValuePtr>{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<Value *>{}); + builder.getSymbolRefAttr(cuGetStreamHelper), ArrayRef<ValuePtr>{}); // Invoke the function with required arguments. auto cuLaunchKernel = getModule().lookupSymbol<LLVM::LLVMFuncOp>(cuLaunchKernelName); @@ -408,19 +408,19 @@ void GpuLaunchFuncToCudaCallsPass::translateGpuLaunchCalls( builder.create<LLVM::CallOp>( loc, ArrayRef<Type>{getCUResultType()}, builder.getSymbolRefAttr(cuLaunchKernel), - 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 */}); + 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 */}); // 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<Value *>(cuStream.getResult(0))); + ArrayRef<ValuePtr>(cuStream.getResult(0))); launchOp.erase(); } |
