[Mlir-commits] [mlir] Revert "[NVGPU] Fix nvdsl examples" (PR #166943)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Fri Nov 7 06:37:19 PST 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir-gpu
Author: Mehdi Amini (joker-eph)
<details>
<summary>Changes</summary>
Reverts llvm/llvm-project#<!-- -->156830
This broke the bots.
---
Patch is 59.24 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/166943.diff
9 Files Affected:
- (modified) mlir/test/Examples/NVGPU/Ch0.py (+2-22)
- (modified) mlir/test/Examples/NVGPU/Ch1.py (+6-42)
- (modified) mlir/test/Examples/NVGPU/Ch2.py (+6-76)
- (modified) mlir/test/Examples/NVGPU/Ch3.py (+4-90)
- (modified) mlir/test/Examples/NVGPU/Ch4.py (+7-187)
- (modified) mlir/test/Examples/NVGPU/Ch5.py (+8-168)
- (modified) mlir/test/Examples/NVGPU/lit.local.cfg (+1-1)
- (modified) mlir/test/Examples/NVGPU/tools/nvdsl.py (+12-14)
- (modified) mlir/test/Examples/NVGPU/tools/nvgpucompiler.py (+1-3)
``````````diff
diff --git a/mlir/test/Examples/NVGPU/Ch0.py b/mlir/test/Examples/NVGPU/Ch0.py
index 6d8ff738feac2..8f60088178d11 100644
--- a/mlir/test/Examples/NVGPU/Ch0.py
+++ b/mlir/test/Examples/NVGPU/Ch0.py
@@ -1,9 +1,5 @@
# RUN: env SUPPORT_LIB=%mlir_cuda_runtime \
-# RUN: sh -c 'if [[ "%mlir_run_cuda_sm90_tests" == "1" ]]; \
-# RUN: then %PYTHON %s | FileCheck %s; \
-# RUN: else export MLIR_NVDSL_PRINT_IR=1; \
-# RUN: %PYTHON %s | FileCheck %s --check-prefix=DUMPIR; fi'
-
+# RUN: %PYTHON %s | FileCheck %s
# ===----------------------------------------------------------------------===//
# Chapter 0 : Hello World
@@ -47,24 +43,8 @@ def kernel():
# 4. The `mlir_func` decorator JIT compiles the IR and executes the MLIR function.
main(alpha)
+
# CHECK: GPU thread 0 has 100
# CHECK: GPU thread 1 has 101
# CHECK: GPU thread 2 has 102
# CHECK: GPU thread 3 has 103
-
-# DUMPIR: func.func @main(%arg0: index) attributes {llvm.emit_c_interface} {
-# DUMPIR: %[[C1:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C1_0:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C1_1:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C4:.*]] = arith.constant 4 : index
-# DUMPIR: %[[C1_2:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C1_3:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C0_I32:.*]] = arith.constant 0 : i32
-# DUMPIR: gpu.launch blocks(%arg1, %arg2, %arg3) in (%arg7 = %[[C1]], %arg8 = %[[C1_0]], %arg9 = %[[C1_1]]) threads(%arg4, %arg5, %arg6) in (%arg10 = %[[C4]], %arg11 = %[[C1_2]], %arg12 = %[[C1_3]]) dynamic_shared_memory_size %[[C0_I32]] {
-# DUMPIR: %[[TIDX:.*]] = gpu.thread_id x
-# DUMPIR: %[[MYVAL:.*]] = arith.addi %arg0, %[[TIDX]] : index
-# DUMPIR: gpu.printf "GPU thread %llu has %llu\0A", %[[TIDX]], %[[MYVAL]] : index, index
-# DUMPIR: gpu.terminator
-# DUMPIR: }
-# DUMPIR: return
-# DUMPIR: }
diff --git a/mlir/test/Examples/NVGPU/Ch1.py b/mlir/test/Examples/NVGPU/Ch1.py
index 0ad740557e8da..cfb48d56f8d49 100644
--- a/mlir/test/Examples/NVGPU/Ch1.py
+++ b/mlir/test/Examples/NVGPU/Ch1.py
@@ -1,9 +1,5 @@
# RUN: env SUPPORT_LIB=%mlir_cuda_runtime \
-# RUN: sh -c 'if [[ "%mlir_run_cuda_sm90_tests" == "1" ]]; \
-# RUN: then %PYTHON %s | FileCheck %s; \
-# RUN: else export MLIR_NVDSL_PRINT_IR=1; \
-# RUN: %PYTHON %s | FileCheck %s --check-prefix=DUMPIR; fi'
-
+# RUN: %PYTHON %s | FileCheck %s
# ===----------------------------------------------------------------------===//
# Chapter 1 : 2D Saxpy
@@ -60,43 +56,11 @@ def saxpy_kernel():
alpha = 2.0
x = np.random.randn(M, N).astype(np.float32)
y = np.ones((M, N), np.float32)
-
saxpy(x, y, alpha)
-if os.getenv("MLIR_NVDSL_PRINT_IR") != "1":
- # 4. Verify MLIR with reference computation
- ref = np.ones((M, N), np.float32)
- ref += x * alpha
- np.testing.assert_allclose(y, ref, rtol=5e-03, atol=1e-01)
- print("PASS")
+# 4. Verify MLIR with reference computation
+ref = np.ones((M, N), np.float32)
+ref += x * alpha
+np.testing.assert_allclose(y, ref, rtol=5e-03, atol=1e-01)
+print("PASS")
# CHECK-NOT: Mismatched elements
-# CHECK: PASS
-
-# DUMPIR: func.func @saxpy(%arg0: memref<256x32xf32>, %arg1: memref<256x32xf32>, %arg2: f32) attributes {llvm.emit_c_interface} {
-# DUMPIR: %[[WAIT0:.*]] = gpu.wait async
-# DUMPIR: %[[MEMREF:.*]], %[[ASYNC0:.*]] = gpu.alloc async [%[[WAIT0]]] () : memref<256x32xf32>
-# DUMPIR: %[[MEMREF0:.*]], %[[ASYNC1:.*]] = gpu.alloc async [%[[ASYNC0]]] () : memref<256x32xf32>
-# DUMPIR: %[[MEMCPY1:.*]] = gpu.memcpy async [%[[ASYNC1]]] %[[MEMREF]], %arg0 : memref<256x32xf32>, memref<256x32xf32>
-# DUMPIR: %[[MEMCPY2:.*]] = gpu.memcpy async [%[[MEMCPY1]]] %[[MEMREF0]], %arg1 : memref<256x32xf32>, memref<256x32xf32>
-# DUMPIR: %[[WAIT1:.*]] = gpu.wait async [%[[MEMCPY2]]]
-# DUMPIR: %[[C256:.*]] = arith.constant 256 : index
-# DUMPIR: %[[C1:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C1_2:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C32:.*]] = arith.constant 32 : index
-# DUMPIR: %[[C1_3:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C1_4:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C0_I32:.*]] = arith.constant 0 : i32
-# DUMPIR: gpu.launch blocks(%arg3, %arg4, %arg5) in (%arg9 = %[[C256]], %arg10 = %[[C1]], %arg11 = %[[C1_2]]) threads(%arg6, %arg7, %arg8) in (%arg12 = %[[C32]], %arg13 = %[[C1_3]], %arg14 = %[[C1_4]]) dynamic_shared_memory_size %[[C0_I32]] {
-# DUMPIR: %[[BLOCKID:.*]] = gpu.block_id x
-# DUMPIR: %[[THREADID:.*]] = gpu.thread_id x
-# DUMPIR: %[[LD0:.*]] = memref.load %[[MEMREF]][%[[BLOCKID]], %[[THREADID]]] : memref<256x32xf32>
-# DUMPIR: %[[LD1:.*]] = memref.load %[[MEMREF0]][%[[BLOCKID]], %[[THREADID]]] : memref<256x32xf32>
-# DUMPIR: %[[MUL:.*]] = arith.mulf %[[LD0]], %arg2 : f32
-# DUMPIR: %[[ADD:.*]] = arith.addf %[[LD1]], %[[MUL]] : f32
-# DUMPIR: memref.store %[[ADD]], %[[MEMREF0]][%[[BLOCKID]], %[[THREADID]]] : memref<256x32xf32>
-# DUMPIR: gpu.terminator
-# DUMPIR: }
-# DUMPIR: %[[MEMCPY3:.*]] = gpu.memcpy async [%[[WAIT1]]] %arg1, %[[MEMREF0]] : memref<256x32xf32>, memref<256x32xf32>
-# DUMPIR: %[[WAIT2:.*]] = gpu.wait async [%[[MEMCPY3]]]
-# DUMPIR: return
-# DUMPIR: }
diff --git a/mlir/test/Examples/NVGPU/Ch2.py b/mlir/test/Examples/NVGPU/Ch2.py
index 18e510221ac42..729913c6d5c4f 100644
--- a/mlir/test/Examples/NVGPU/Ch2.py
+++ b/mlir/test/Examples/NVGPU/Ch2.py
@@ -1,9 +1,5 @@
# RUN: env SUPPORT_LIB=%mlir_cuda_runtime \
-# RUN: sh -c 'if [[ "%mlir_run_cuda_sm90_tests" == "1" ]]; \
-# RUN: then %PYTHON %s | FileCheck %s; \
-# RUN: else export MLIR_NVDSL_PRINT_IR=1; \
-# RUN: %PYTHON %s | FileCheck %s --check-prefix=DUMPIR; fi'
-
+# RUN: %PYTHON %s | FileCheck %s
# ===----------------------------------------------------------------------===//
# Chapter 2 : 2D Saxpy with TMA
@@ -89,75 +85,9 @@ def saxpy_tma_kernel():
y = np.ones((M, N), np.float32)
saxpy(x, y, alpha)
-if os.getenv("MLIR_NVDSL_PRINT_IR") != "1":
- # 4. Verify MLIR with reference computation
- ref = np.ones((M, N), np.float32)
- ref += x * alpha
- np.testing.assert_allclose(y, ref, rtol=5e-03, atol=1e-01)
- print("PASS")
+# 4. Verify MLIR with reference computation
+ref = np.ones((M, N), np.float32)
+ref += x * alpha
+np.testing.assert_allclose(y, ref, rtol=5e-03, atol=1e-01)
+print("PASS")
# CHECK-NOT: Mismatched elements
-# CHECK: PASS
-
-# DUMPIR: func.func @saxpy(%arg0: memref<256x32xf32>, %arg1: memref<256x32xf32>, %arg2: f32) attributes {llvm.emit_c_interface} {
-# DUMPIR: %[[WAIT0:.*]] = gpu.wait async
-# DUMPIR: %[[MEMREF:.*]], %[[ASYNC0:.*]] = gpu.alloc async [%[[WAIT0]]] () : memref<256x32xf32>
-# DUMPIR: %[[MEMREF0:.*]], %[[ASYNC1:.*]] = gpu.alloc async [%[[ASYNC0]]] () : memref<256x32xf32>
-# DUMPIR: %[[MEMCPY1:.*]] = gpu.memcpy async [%[[ASYNC1]]] %[[MEMREF]], %arg0 : memref<256x32xf32>, memref<256x32xf32>
-# DUMPIR: %[[MEMCPY2:.*]] = gpu.memcpy async [%[[MEMCPY1]]] %[[MEMREF0]], %arg1 : memref<256x32xf32>, memref<256x32xf32>
-# DUMPIR: %[[WAIT1:.*]] = gpu.wait async [%[[MEMCPY2]]]
-# DUMPIR: %[[CAST:.*]] = memref.cast %[[MEMREF]] : memref<256x32xf32> to memref<*xf32>
-# DUMPIR: %[[C1:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C32:.*]] = arith.constant 32 : index
-# DUMPIR: %[[TMA0:.*]] = nvgpu.tma.create.descriptor %[[CAST]] box[%[[C1]], %[[C32]]] : memref<*xf32> -> <tensor = memref<1x32xf32, 3>, swizzle = none, l2promo = none, oob = zero, interleave = none>
-# DUMPIR: %[[CAST2:.*]] = memref.cast %[[MEMREF0]] : memref<256x32xf32> to memref<*xf32>
-# DUMPIR: %[[C1_3:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C32_4:.*]] = arith.constant 32 : index
-# DUMPIR: %[[TMA1:.*]] = nvgpu.tma.create.descriptor %[[CAST2]] box[%[[C1_3]], %[[C32_4]]] : memref<*xf32> -> <tensor = memref<1x32xf32, 3>, swizzle = none, l2promo = none, oob = zero, interleave = none>
-# DUMPIR: %[[C256:.*]] = arith.constant 256 : index
-# DUMPIR: %[[C1_5:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C1_6:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C32_7:.*]] = arith.constant 32 : index
-# DUMPIR: %[[C1_8:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C1_9:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C256_I32:.*]] = arith.constant 256 : i32
-# DUMPIR: gpu.launch blocks(%arg3, %arg4, %arg5) in (%arg9 = %[[C256]], %arg10 = %[[C1_5]], %arg11 = %[[C1_6]]) threads(%arg6, %arg7, %arg8) in (%arg12 = %[[C32_7]], %arg13 = %[[C1_8]], %arg14 = %[[C1_9]]) dynamic_shared_memory_size %[[C256_I32]] {
-# DUMPIR: %[[BLOCKID:.*]] = gpu.block_id x
-# DUMPIR: %[[THREADID:.*]] = gpu.thread_id x
-# DUMPIR: %[[C0:.*]] = arith.constant 0 : index
-# DUMPIR: %[[EQ:.*]] = arith.cmpi eq, %[[THREADID]], %[[C0]] : index
-# DUMPIR: %[[MB:.*]] = nvgpu.mbarrier.create -> <memorySpace = #gpu.address_space<workgroup>>
-# DUMPIR: %[[C0_10:.*]] = arith.constant 0 : index
-# DUMPIR: %[[C1_11:.*]] = arith.constant 1 : index
-# DUMPIR: nvgpu.mbarrier.init %[[MB]][%[[C0_10]]], %[[C1_11]], predicate = %[[EQ]] : <memorySpace = #gpu.address_space<workgroup>>
-# DUMPIR: %[[DSM0:.*]] = gpu.dynamic_shared_memory : memref<?xi8, #gpu.address_space<workgroup>>
-# DUMPIR: %[[C0_12:.*]] = arith.constant 0 : index
-# DUMPIR: %[[VIEW:.*]] = memref.view %[[DSM0]][%[[C0_12]]][] : memref<?xi8, #gpu.address_space<workgroup>> to memref<1x32xf32, #gpu.address_space<workgroup>>
-# DUMPIR: %[[DSM1:.*]] = gpu.dynamic_shared_memory : memref<?xi8, #gpu.address_space<workgroup>>
-# DUMPIR: %[[C128:.*]] = arith.constant 128 : index
-# DUMPIR: %[[VIEW_13:.*]] = memref.view %[[DSM1]][%[[C128]]][] : memref<?xi8, #gpu.address_space<workgroup>> to memref<1x32xf32, #gpu.address_space<workgroup>>
-# DUMPIR: %[[C0_14:.*]] = arith.constant 0 : index
-# DUMPIR: %[[C0_15:.*]] = arith.constant 0 : index
-# DUMPIR: nvgpu.tma.async.load %[[TMA0]][%[[C0_15]], %[[BLOCKID]]], %[[MB]][%[[C0_14]]] to %[[VIEW]], predicate = %[[EQ]] : <tensor = memref<1x32xf32, 3>, swizzle = none, l2promo = none, oob = zero, interleave = none>, <memorySpace = #gpu.address_space<workgroup>> -> memref<1x32xf32, #gpu.address_space<workgroup>>
-# DUMPIR: %[[C0_16:.*]] = arith.constant 0 : index
-# DUMPIR: %[[C0_17:.*]] = arith.constant 0 : index
-# DUMPIR: nvgpu.tma.async.load %[[TMA1]][%[[C0_17]], %[[BLOCKID]]], %[[MB]][%[[C0_16]]] to %[[VIEW_13]], predicate = %[[EQ]] : <tensor = memref<1x32xf32, 3>, swizzle = none, l2promo = none, oob = zero, interleave = none>, <memorySpace = #gpu.address_space<workgroup>> -> memref<1x32xf32, #gpu.address_space<workgroup>>
-# DUMPIR: %[[C0_18:.*]] = arith.constant 0 : index
-# DUMPIR: %[[C256_19:.*]] = arith.constant 256 : index
-# DUMPIR: nvgpu.mbarrier.arrive.expect_tx %[[MB]][%[[C0_18]]], %[[C256_19]], predicate = %[[EQ]] : <memorySpace = #gpu.address_space<workgroup>>
-# DUMPIR: %[[C0_20:.*]] = arith.constant 0 : index
-# DUMPIR: %[[C10000000:.*]] = arith.constant 10000000 : index
-# DUMPIR: %[[FALSE:.*]] = arith.constant false
-# DUMPIR: nvgpu.mbarrier.try_wait.parity %[[MB]][%[[C0_20]]], %[[FALSE]], %[[C10000000]] : <memorySpace = #gpu.address_space<workgroup>>
-# DUMPIR: %[[C0_21:.*]] = arith.constant 0 : index
-# DUMPIR: %[[LD0:.*]] = memref.load %[[VIEW]][%[[C0_21]], %[[THREADID]]] : memref<1x32xf32, #gpu.address_space<workgroup>>
-# DUMPIR: %[[C0_22:.*]] = arith.constant 0 : index
-# DUMPIR: %[[LD1:.*]] = memref.load %[[VIEW_13]][%[[C0_22]], %[[THREADID]]] : memref<1x32xf32, #gpu.address_space<workgroup>>
-# DUMPIR: %[[MUL:.*]] = arith.mulf %[[LD0]], %arg2 : f32
-# DUMPIR: %[[ADD:.*]] = arith.addf %[[LD1]], %[[MUL]] : f32
-# DUMPIR: memref.store %[[ADD]], %[[MEMREF0]][%[[BLOCKID]], %[[THREADID]]] : memref<256x32xf32>
-# DUMPIR: gpu.terminator
-# DUMPIR: }
-# DUMPIR: %[[MEMCPY3:.*]] = gpu.memcpy async [%[[WAIT1]]] %arg1, %[[MEMREF0]] : memref<256x32xf32>, memref<256x32xf32>
-# DUMPIR: %[[WAIT2:.*]] = gpu.wait async [%[[MEMCPY3]]]
-# DUMPIR: return
-# DUMPIR: }
diff --git a/mlir/test/Examples/NVGPU/Ch3.py b/mlir/test/Examples/NVGPU/Ch3.py
index dbc73c3efc60c..eb96b11c63416 100644
--- a/mlir/test/Examples/NVGPU/Ch3.py
+++ b/mlir/test/Examples/NVGPU/Ch3.py
@@ -1,9 +1,5 @@
# RUN: env SUPPORT_LIB=%mlir_cuda_runtime \
-# RUN: sh -c 'if [[ "%mlir_run_cuda_sm90_tests" == "1" ]]; \
-# RUN: then %PYTHON %s | FileCheck %s; \
-# RUN: else export MLIR_NVDSL_PRINT_IR=1; \
-# RUN: %PYTHON %s | FileCheck %s --check-prefix=DUMPIR; fi'
-
+# RUN: %PYTHON %s | FileCheck %s
# ===----------------------------------------------------------------------===//
# Chapter 3 : GEMM 128x128x64 with Tensor Core
@@ -127,89 +123,7 @@ def gemm_tma_kernel():
d = np.zeros((M, N), np.float32)
gemm_128_128_64(a, b, d)
-if os.getenv("MLIR_NVDSL_PRINT_IR") != "1":
- # Verify MLIR program with reference computation in python
- ref_d = a.astype(np.float16) @ b.astype(np.float16)
- np.testing.assert_allclose(d, ref_d, rtol=5e-03, atol=1e-01)
- print("PASS")
+ref_d = a.astype(np.float16) @ b.astype(np.float16)
+np.testing.assert_allclose(d, ref_d, rtol=5e-03, atol=1e-01)
+print("PASS")
# CHECK-NOT: Mismatched elements
-# CHECK: PASS
-
-# DUMPIR: func.func @gemm_128_128_64(%arg0: memref<128x64xf16>, %arg1: memref<64x128xf16>, %arg2: memref<128x128xf32>) attributes {llvm.emit_c_interface} {
-# DUMPIR: %[[WAIT0:.*]] = gpu.wait async
-# DUMPIR: %[[MEM0:.*]], %[[ASYNC0:.*]] = gpu.alloc async [%[[WAIT0]]] () : memref<128x64xf16>
-# DUMPIR: %[[MEM1:.*]], %[[ASYNC1:.*]] = gpu.alloc async [%[[ASYNC0]]] () : memref<64x128xf16>
-# DUMPIR: %[[MEM2:.*]], %[[ASYNC2:.*]] = gpu.alloc async [%[[ASYNC1]]] () : memref<128x128xf32>
-# DUMPIR: %[[CPY1:.*]] = gpu.memcpy async [%[[ASYNC2]]] %[[MEM0]], %arg0 : memref<128x64xf16>, memref<128x64xf16>
-# DUMPIR: %[[CPY2:.*]] = gpu.memcpy async [%[[CPY1]]] %[[MEM1]], %arg1 : memref<64x128xf16>, memref<64x128xf16>
-# DUMPIR: %[[WAIT1:.*]] = gpu.wait async [%[[CPY2]]]
-# DUMPIR: %[[CAST0:.*]] = memref.cast %[[MEM0]] : memref<128x64xf16> to memref<*xf16>
-# DUMPIR: %[[C128:.*]] = arith.constant 128 : index
-# DUMPIR: %[[C64:.*]] = arith.constant 64 : index
-# DUMPIR: %[[TMA0:.*]] = nvgpu.tma.create.descriptor %[[CAST0]] box[%[[C128]], %[[C64]]] : memref<*xf16> -> <tensor = memref<128x64xf16, 3>, swizzle = swizzle_128b, l2promo = none, oob = zero, interleave = none>
-# DUMPIR: %[[CAST1:.*]] = memref.cast %[[MEM1]] : memref<64x128xf16> to memref<*xf16>
-# DUMPIR: %[[C64_5:.*]] = arith.constant 64 : index
-# DUMPIR: %[[C64_6:.*]] = arith.constant 64 : index
-# DUMPIR: %[[TMA1:.*]] = nvgpu.tma.create.descriptor %[[CAST1]] box[%[[C64_5]], %[[C64_6]]] : memref<*xf16> -> <tensor = memref<64x64xf16, 3>, swizzle = swizzle_128b, l2promo = none, oob = zero, interleave = none>
-# DUMPIR: %[[C1:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C1_7:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C1_8:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C128_9:.*]] = arith.constant 128 : index
-# DUMPIR: %[[C1_10:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C1_11:.*]] = arith.constant 1 : index
-# DUMPIR: %[[C32768_I32:.*]] = arith.constant 32768 : i32
-# DUMPIR: gpu.launch blocks(%arg3, %arg4, %arg5) in (%arg9 = %[[C1]], %arg10 = %[[C1_7]], %arg11 = %[[C1_8]]) threads(%arg6, %arg7, %arg8) in (%arg12 = %[[C128_9]], %arg13 = %[[C1_10]], %arg14 = %[[C1_11]]) dynamic_shared_memory_size %[[C32768_I32]] {
-# DUMPIR: %[[THREADID:.*]] = gpu.thread_id x
-# DUMPIR: %[[MB:.*]] = nvgpu.mbarrier.create -> <memorySpace = #gpu.address_space<workgroup>>
-# DUMPIR: %[[C0:.*]] = arith.constant 0 : index
-# DUMPIR: %[[EQ:.*]] = arith.cmpi eq, %[[THREADID]], %[[C0]] : index
-# DUMPIR: %[[C0_12:.*]] = arith.constant 0 : index
-# DUMPIR: %[[C1_13:.*]] = arith.constant 1 : index
-# DUMPIR: nvgpu.mbarrier.init %[[MB]][%[[C0_12]]], %[[C1_13]], predicate = %[[EQ]] : <memorySpace = #gpu.address_space<workgroup>>
-# DUMPIR: nvgpu.tma.prefetch.descriptor %[[TMA0]], predicate = %[[EQ]] : <tensor = memref<128x64xf16, 3>, swizzle = swizzle_128b, l2promo = none, oob = zero, interleave = none>
-# DUMPIR: nvgpu.tma.prefetch.descriptor %[[TMA1]], predicate = %[[EQ]] : <tensor = memref<64x64xf16, 3>, swizzle = swizzle_128b, l2promo = none, oob = zero, interleave = none>
-# DUMPIR: %[[DSM0:.*]] = gpu.dynamic_shared_memory : memref<?xi8, #gpu.address_space<workgroup>>
-# DUMPIR: %[[C0_14:.*]] = arith.constant 0 : index
-# DUMPIR: %[[VIEW:.*]] = memref.view %[[DSM0]][%[[C0_14]]][] : memref<?xi8, #gpu.address_space<workgroup>> to memref<128x64xf16, #gpu.address_space<workgroup>>
-# DUMPIR: %[[DSM1:.*]] = gpu.dynamic_shared_memory : memref<?xi8, #gpu.address_space<workgroup>>
-# DUMPIR: %[[C16384:.*]] = arith.constant 16384 : index
-# DUMPIR: %[[VIEW_15:.*]] = memref.view %[[DSM1]][%[[C16384]]][] : memref<?xi8, #gpu.address_space<workgroup>> to memref<64x128xf16, #gpu.address_space<workgroup>>
-# DUMPIR: %[[DSM2:.*]] = gpu.dynamic_shared_memory : memref<?xi8, #gpu.address_space<workgroup>>
-# DUMPIR: %[[C0_16:.*]] = arith.constant 0 : index
-# DUMPIR: %[[VIEW_17:.*]] = memref.view %[[DSM2]][%[[C0_16]]][] : memref<?xi8, #gpu.address_space<workgroup>> to memref<128x64xf16, #gpu.address_space<workgroup>>
-# DUMPIR: %[[DSM3:.*]] = gpu.dynamic_shared_memory : memref<?xi8, #gpu.address_space<workgroup>>
-# DUMPIR: %[[C16384_18:.*]] = arith.constant 16384 : index
-# DUMPIR: %[[VIEW_19:.*]] = memref.view %[[DSM3]][%[[C16384_18]]][] : memref<?xi8, #gpu.address_space<workgroup>> to memref<64x64xf16, #gpu.address_space<workgroup>>
-# DUMPIR: %[[DSM4:.*]] = gpu.dynamic_shared_memory : memref<?xi8, #gpu.address_space<workgroup>>
-# DUMPIR: %[[C24576:.*]] = arith.constant 24576 : index
-# DUMPIR: %[[VIEW_20:.*]] = memref.view %[[DSM4]][%[[C24576]]][] : memref<?xi8, #gpu.address_space<workgroup>> to memref<64x64xf16, #gpu.address_space<workgroup>>
-# DUMPIR: %[[C0_21:.*]] = arith.constant 0 : index
-# DUMPIR: %[[C32768:.*]] = arith.constant 32768 : index
-# DUMPIR: nvgpu.mbarrier.arrive.expect_tx %[[MB]][%[[C0_21]]], %[[C32768]], predicate = %[[EQ]] : <memorySpace = #gpu.address_space<workgroup>>
-# DUMPIR: %[[C0_22:.*]] = arith.constant 0 : index
-# DUMPIR: %[[C0_23:.*]] = arith.constant 0 : index
-# DUMPIR: %[[C0_24:.*]] = arith.constant 0 : index
-# DUMPIR: nvgpu.tma.async.load %[[TMA0]][%[[C0_23]], %[[C0_24]]], %[[MB]][%[[C0_22]]] to %[[VIEW_17]], predicate = %[[EQ]] : <tensor = memref<128x64xf16, 3>, swizzle = swizzle_128b, l2promo = none, oob = zero, interleave = none>, <memorySpace = #gpu.address_space<workgroup>> -> memref<128x64xf16, #gpu.address_space<workgroup>>
-# DUMPIR: %[[C0_25:.*]] = arith.constant 0 : index
-# DUMPIR: %[[C0_26:.*]] = arith.constant 0 : index
-# DUMPIR: %[[C0_27:.*]] = arith.constant 0 : index
-# DUMPIR: nvgpu.tma.async.load %[[TMA1]][%[[C0_26]], %[[C0_27]]], %[[MB]][%[[C0_25]]] to %[[VIEW_19]], predicate = %[[EQ]] : <tensor = memref<64x64xf16, 3>, swizzle = swizzle_128b, l2promo = none, oob = zero, interleave = none>, <memorySpace = #gpu.address_space<workgroup>> -> memref<64x64xf16, #gpu.address_space<workgroup>>
-# DUMPIR: %[[C0_28:.*]] = arith.constant 0 : index
-# DUMPIR: %[[C64_29:.*]] = arith.constant 64 : index
-# DUMPIR: %[[C0_30:.*]] = arith.constant 0 : index
-# DUMPIR: nvgpu.tma.async.load %[[TMA1]][%[[C64_29]], %[[C0_30]]], %[[MB]][%[[C0_28]]] to %[[VIEW_20]], predicate = %[[EQ]] : <tensor = memref<64x64xf16, 3>, swizzle = swizzle_128b, l2promo = none, oob = zero, interleave = none>, <memorySpace = #gpu.address_space<workgroup>> -> memref<64x64xf16, #gpu.address_space<workgroup>>
-# DUMPIR: %[[C0_31:.*]] = arith.constant 0 : index
-# DUMPIR: %[[C10000000:.*]] = arith.constant 10000000 : index
-# DUMPIR: %[[FALSE:.*]] = arith.constant false
-# DUMPIR: nvgpu.mbarrier.try_wait.parity %[[MB]][%[[C0_31]]], %[[FALSE]], %[[C10000000]] : <memorySpace = #gpu.address_space<workgroup>>
-# DUMPIR: %[[WG_ACC:.*]] = nvgpu.warpgroup.mma.init.accumulator -> <...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/166943
More information about the Mlir-commits
mailing list