intel/llvm
1
0
Fork 0
mirror of https://github.com/intel/llvm.git synced 2026-01-26 12:26:52 +08:00
Code Issues Packages Projects Releases Wiki Activity

[MLIR][GPUToROCDL] Remove typed pointer support (#70908)

Browse Source 
This commit removes the support for lowering GPU to ROCDL dialect with
typed pointers. Typed pointers have been deprecated for a while now and
it's planned to soon remove them from the LLVM dialect.

Related PSA:
https://discourse.llvm.org/t/psa-removal-of-typed-pointers-from-the-llvm-dialect/74502
This commit is contained in:
Christian Ulmann
2023-11-01 10:13:06 +01:00
committed by GitHub
parent b120fe8d32
commit 4279a642fb
12 changed files with 12 additions and 50 deletions
Download Patch File Download Diff File Expand all files Collapse all files

5
mlir/include/mlir/Conversion/Passes.td
View File

@@ -529,10 +529,7 @@ def ConvertGpuOpsToROCDLOps : Pass<"convert-gpu-to-rocdl", "gpu::GPUModuleOp"> {
clEnumValN(::mlir::gpu::amd::Runtime::Unknown, "unknown", "Unknown (default)"),
clEnumValN(::mlir::gpu::amd::Runtime::HIP, "HIP", "HIP"),
clEnumValN(::mlir::gpu::amd::Runtime::OpenCL, "OpenCL", "OpenCL")
)}]>,
Option<"useOpaquePointers", "use-opaque-pointers", "bool",
/*default=*/"true", "Generate LLVM IR using opaque pointers "
"instead of typed pointers">,
)}]>
];
}

1
mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
View File

@@ -229,7 +229,6 @@ struct LowerGpuOpsToROCDLOpsPass
ctx, DataLayout(cast<DataLayoutOpInterface>(m.getOperation())));
if (indexBitwidth != kDeriveIndexBitwidthFromDataLayout)
options.overrideIndexBitwidth(indexBitwidth);
options.useOpaquePointers = useOpaquePointers;
if (useBarePtrCallConv) {
options.useBarePtrCallConv = true;

2
mlir/test/Conversion/GPUCommon/lower-memory-space-attrs.mlir
View File

@@ -1,4 +1,4 @@
// RUN: mlir-opt %s -split-input-file -convert-gpu-to-rocdl='use-opaque-pointers=1' | FileCheck %s --check-prefixes=CHECK,ROCDL
// RUN: mlir-opt %s -split-input-file -convert-gpu-to-rocdl | FileCheck %s --check-prefixes=CHECK,ROCDL
// RUN: mlir-opt %s -split-input-file -convert-gpu-to-nvvm | FileCheck %s --check-prefixes=CHECK,NVVM
gpu.module @kernel {

2
mlir/test/Conversion/GPUCommon/memory-attrbution.mlir
View File

@@ -1,5 +1,5 @@
// RUN: mlir-opt -allow-unregistered-dialect --convert-gpu-to-nvvm --split-input-file %s | FileCheck --check-prefix=NVVM %s
// RUN: mlir-opt -allow-unregistered-dialect --convert-gpu-to-rocdl='use-opaque-pointers=1' --split-input-file %s | FileCheck --check-prefix=ROCDL %s
// RUN: mlir-opt -allow-unregistered-dialect --convert-gpu-to-rocdl --split-input-file %s | FileCheck --check-prefix=ROCDL %s
gpu.module @kernel {
// NVVM-LABEL: llvm.func @private

2
mlir/test/Conversion/GPUCommon/memref-arg-attrs.mlir
View File

@@ -1,4 +1,4 @@
// RUN: mlir-opt %s -split-input-file -convert-gpu-to-rocdl='use-opaque-pointers=1 use-bare-ptr-memref-call-conv=0' | FileCheck %s --check-prefixes=CHECK,ROCDL
// RUN: mlir-opt %s -split-input-file -convert-gpu-to-rocdl='use-bare-ptr-memref-call-conv=0' | FileCheck %s --check-prefixes=CHECK,ROCDL
// RUN: mlir-opt %s -split-input-file -convert-gpu-to-nvvm='use-bare-ptr-memref-call-conv=0' | FileCheck %s --check-prefixes=CHECK,NVVM
gpu.module @kernel {

2
mlir/test/Conversion/GPUCommon/memref-arg-noalias-attrs.mlir
View File

@@ -1,4 +1,4 @@
// RUN: mlir-opt %s -split-input-file -convert-gpu-to-rocdl='use-opaque-pointers=1 use-bare-ptr-memref-call-conv=1' | FileCheck %s --check-prefixes=CHECK,ROCDL
// RUN: mlir-opt %s -split-input-file -convert-gpu-to-rocdl='use-bare-ptr-memref-call-conv=1' | FileCheck %s --check-prefixes=CHECK,ROCDL
// RUN: mlir-opt %s -split-input-file -convert-gpu-to-nvvm='use-bare-ptr-memref-call-conv=1' | FileCheck %s --check-prefixes=CHECK,NVVM
gpu.module @kernel {

2
mlir/test/Conversion/GPUCommon/memref-arg-noalias-warning.mlir
View File

@@ -1,4 +1,4 @@
// RUN: mlir-opt %s -split-input-file -convert-gpu-to-rocdl='use-opaque-pointers=1 use-bare-ptr-memref-call-conv=0' -verify-diagnostics
// RUN: mlir-opt %s -split-input-file -convert-gpu-to-rocdl='use-bare-ptr-memref-call-conv=0' -verify-diagnostics
gpu.module @kernel {
// expected-warning @+1 {{Cannot copy noalias with non-bare pointers.}}

2
mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-hip.mlir
View File

@@ -1,4 +1,4 @@
// RUN: mlir-opt %s -convert-gpu-to-rocdl='runtime=HIP use-opaque-pointers=1' -split-input-file | FileCheck %s
// RUN: mlir-opt %s -convert-gpu-to-rocdl='runtime=HIP' -split-input-file | FileCheck %s
gpu.module @test_module {
// CHECK-DAG: llvm.mlir.global internal constant @[[$PRINT_GLOBAL0:[A-Za-z0-9_]+]]("Hello, world\0A\00")

2
mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-opencl.mlir
View File

@@ -1,4 +1,4 @@
// RUN: mlir-opt %s -convert-gpu-to-rocdl='runtime=OpenCL use-opaque-pointers=1' | FileCheck %s
// RUN: mlir-opt %s -convert-gpu-to-rocdl='runtime=OpenCL' | FileCheck %s
gpu.module @test_module {
// CHECK: llvm.mlir.global internal constant @[[$PRINT_GLOBAL:[A-Za-z0-9_]+]]("Hello: %d\0A\00") {addr_space = 4 : i32}

4
mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir
View File

@@ -1,5 +1,5 @@
// RUN: mlir-opt %s -convert-gpu-to-rocdl='use-opaque-pointers=1' -split-input-file | FileCheck %s
// RUN: mlir-opt %s -convert-gpu-to-rocdl='index-bitwidth=32 use-opaque-pointers=1' -split-input-file | FileCheck --check-prefix=CHECK32 %s
// RUN: mlir-opt %s -convert-gpu-to-rocdl -split-input-file | FileCheck %s
// RUN: mlir-opt %s -convert-gpu-to-rocdl='index-bitwidth=32' -split-input-file | FileCheck --check-prefix=CHECK32 %s
gpu.module @test_module {
// CHECK-LABEL: func @gpu_index_ops()

4
mlir/test/Conversion/GPUToROCDL/memref.mlir
View File

@@ -1,6 +1,6 @@
// RUN: mlir-opt %s -convert-gpu-to-rocdl='use-opaque-pointers=1' -split-input-file | FileCheck %s
// RUN: mlir-opt %s -convert-gpu-to-rocdl -split-input-file | FileCheck %s
// RUN: mlir-opt %s \
// RUN: -convert-gpu-to-rocdl='use-bare-ptr-memref-call-conv=true use-opaque-pointers=1' \
// RUN: -convert-gpu-to-rocdl='use-bare-ptr-memref-call-conv=true' \
// RUN: -split-input-file \
// RUN: | FileCheck %s --check-prefix=BARE

34
mlir/test/Conversion/GPUToROCDL/typed-pointers.mlir
View File

@@ -1,34 +0,0 @@
// RUN: mlir-opt %s -convert-gpu-to-rocdl="runtime=HIP use-opaque-pointers=0" -split-input-file | FileCheck --check-prefixes=CHECK,HIP %s
// RUN: mlir-opt %s -convert-gpu-to-rocdl="runtime=OpenCL use-opaque-pointers=0" | FileCheck --check-prefixes=CHECK,OCL %s
gpu.module @test_module {
// HIP-DAG: llvm.mlir.global internal constant @[[$PRINT_GLOBAL1:[A-Za-z0-9_]+]]("Hello: %d\0A\00")
// HIP-DAG: llvm.func @__ockl_printf_append_args(i64, i32, i64, i64, i64, i64, i64, i64, i64, i32) -> i64
// HIP-DAG: llvm.func @__ockl_printf_append_string_n(i64, !llvm.ptr<i8>, i64, i32) -> i64
// HIP-DAG: llvm.func @__ockl_printf_begin(i64) -> i64
// OCL: llvm.mlir.global internal constant @[[$PRINT_GLOBAL:[A-Za-z0-9_]+]]("Hello: %d\0A\00") {addr_space = 4 : i32}
// OCL: llvm.func @printf(!llvm.ptr<i8, 4>, ...) -> i32
// CHECK-LABEL: func @test_printf
// CHECK: (%[[ARG0:.*]]: i32)
gpu.func @test_printf(%arg0: i32) {
// OCL: %[[IMM0:.*]] = llvm.mlir.addressof @[[$PRINT_GLOBAL]] : !llvm.ptr<array<11 x i8>, 4>
// OCL-NEXT: %[[IMM2:.*]] = llvm.getelementptr %[[IMM0]][0, 0] : (!llvm.ptr<array<11 x i8>, 4>) -> !llvm.ptr<i8, 4>
// OCL-NEXT: %{{.*}} = llvm.call @printf(%[[IMM2]], %[[ARG0]]) vararg(!llvm.func<i32 (ptr<i8, 4>, ...)>) : (!llvm.ptr<i8, 4>, i32) -> i32
// HIP: %[[CST0:.*]] = llvm.mlir.constant(0 : i64) : i64
// HIP-NEXT: %[[DESC0:.*]] = llvm.call @__ockl_printf_begin(%0) : (i64) -> i64
// HIP-NEXT: %[[FORMATSTR:.*]] = llvm.mlir.addressof @[[$PRINT_GLOBAL1]] : !llvm.ptr<array<11 x i8>>
// HIP-NEXT: %[[FORMATSTART:.*]] = llvm.getelementptr %[[FORMATSTR]][0, 0] : (!llvm.ptr<array<11 x i8>>) -> !llvm.ptr<i8>
// HIP-NEXT: %[[FORMATLEN:.*]] = llvm.mlir.constant(11 : i64) : i64
// HIP-NEXT: %[[ISLAST:.*]] = llvm.mlir.constant(1 : i32) : i32
// HIP-NEXT: %[[ISNTLAST:.*]] = llvm.mlir.constant(0 : i32) : i32
// HIP-NEXT: %[[DESC1:.*]] = llvm.call @__ockl_printf_append_string_n(%[[DESC0]], %[[FORMATSTART]], %[[FORMATLEN]], %[[ISNTLAST]]) : (i64, !llvm.ptr<i8>, i64, i32) -> i64
// HIP-NEXT: %[[NARGS1:.*]] = llvm.mlir.constant(1 : i32) : i32
// HIP-NEXT: %[[ARG0_64:.*]] = llvm.zext %[[ARG0]] : i32 to i64
// HIP-NEXT: %{{.*}} = llvm.call @__ockl_printf_append_args(%[[DESC1]], %[[NARGS1]], %[[ARG0_64]], %[[CST0]], %[[CST0]], %[[CST0]], %[[CST0]], %[[CST0]], %[[CST0]], %[[ISLAST]]) : (i64, i32, i64, i64, i64, i64, i64, i64, i64, i32) -> i64
gpu.printf "Hello: %d\n" %arg0 : i32
gpu.return
}
}