Author: Valentin Clement (バレンタイン クレメン) Date: 2024-12-04T10:16:54-08:00 New Revision: 6abd04e62cb913d970e74b13486a0485e175cf02
URL: https://github.com/llvm/llvm-project/commit/6abd04e62cb913d970e74b13486a0485e175cf02 DIFF: https://github.com/llvm/llvm-project/commit/6abd04e62cb913d970e74b13486a0485e175cf02.diff LOG: Revert "[flang][cuda] Run target rewrite in gpu.module (#118592)" This reverts commit cd92c6a89541cbbb67b39142d93a76caae0f79bf. Added: Modified: flang/lib/Optimizer/CodeGen/TargetRewrite.cpp Removed: flang/test/Fir/CUDA/cuda-target-rewrite.mlir ################################################################################ diff --git a/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp b/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp index 1b86d5241704b1..ae6e7ce798d998 100644 --- a/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp +++ b/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp @@ -27,7 +27,6 @@ #include "flang/Optimizer/Dialect/Support/FIRContext.h" #include "flang/Optimizer/Support/DataLayout.h" #include "mlir/Dialect/DLTI/DLTI.h" -#include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Transforms/DialectConversion.h" #include "llvm/ADT/STLExtras.h" @@ -721,11 +720,6 @@ class TargetRewrite : public fir::impl::TargetRewritePassBase<TargetRewrite> { convertSignature(fn); } - - for (auto gpuMod : mod.getOps<mlir::gpu::GPUModuleOp>()) - for (auto fn : gpuMod.getOps<mlir::func::FuncOp>()) - convertSignature(fn); - return mlir::success(); } diff --git a/flang/test/Fir/CUDA/cuda-target-rewrite.mlir b/flang/test/Fir/CUDA/cuda-target-rewrite.mlir deleted file mode 100644 index d85cca38870adb..00000000000000 --- a/flang/test/Fir/CUDA/cuda-target-rewrite.mlir +++ /dev/null @@ -1,16 +0,0 @@ -// RUN: fir-opt --target-rewrite %s | FileCheck %s - -gpu.module @testmod { - gpu.func @_QPvcpowdk(%arg0: !fir.ref<complex<f64>> {cuf.data_attr = #cuf.cuda<device>, fir.bindc_name = "a"}) attributes {cuf.proc_attr = #cuf.cuda_proc<global>} { - %0 = fir.alloca i64 - %1 = fir.load %0 : !fir.ref<i64> - %2 = fir.load %arg0 : !fir.ref<complex<f64>> - %3 = fir.call @_FortranAzpowk(%2, %1) fastmath<contract> : (complex<f64>, i64) -> complex<f64> - gpu.return - } - func.func private @_FortranAzpowk(complex<f64>, i64) -> complex<f64> attributes {fir.bindc_name = "_FortranAzpowk", fir.runtime} -} - -// CHECK-LABEL: gpu.func @_QPvcpowdk -// CHECK: %{{.*}} = fir.call @_FortranAzpowk(%{{.*}}, %{{.*}}, %{{.*}}) : (f64, f64, i64) -> tuple<f64, f64> -// CHECK: func.func private @_FortranAzpowk(f64, f64, i64) -> tuple<f64, f64> attributes {fir.bindc_name = "_FortranAzpowk", fir.runtime} _______________________________________________ llvm-branch-commits mailing list llvm-branch-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
