diff options
| author | Valentin Clement (バレンタイン クレメン) <clementval@gmail.com> | 2025-11-03 13:11:10 -1000 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2025-11-03 23:11:10 +0000 |
| commit | 73ef4dd48fa3ef64a0f291cb822ab66289780b0b (patch) | |
| tree | 27bd9ae53a8c152d51e37f0062d913a4ff3e1774 | |
| parent | 615299934489953deaf202cc445ac9f8ad362afc (diff) | |
| download | llvm-73ef4dd48fa3ef64a0f291cb822ab66289780b0b.zip llvm-73ef4dd48fa3ef64a0f291cb822ab66289780b0b.tar.gz llvm-73ef4dd48fa3ef64a0f291cb822ab66289780b0b.tar.bz2 | |
[flang][cuda] Add missing semi-colon in inlined ptx (#166254)
This would trigger error in ptxas.
| -rw-r--r-- | flang/lib/Optimizer/Builder/IntrinsicCall.cpp | 4 | ||||
| -rw-r--r-- | flang/test/Lower/CUDA/cuda-device-proc.cuf | 16 |
2 files changed, 10 insertions, 10 deletions
diff --git a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp index 2db0606..6ebd52d 100644 --- a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp +++ b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp @@ -9501,7 +9501,7 @@ void IntrinsicLibrary::genTMABulkS2G(llvm::ArrayRef<fir::ExtendedValue> args) { builder, loc, dst, src, fir::getBase(args[2]), {}, {}); mlir::NVVM::InlinePtxOp::create(builder, loc, mlir::TypeRange{}, {}, {}, - "cp.async.bulk.commit_group", {}); + "cp.async.bulk.commit_group;", {}); mlir::NVVM::CpAsyncBulkWaitGroupOp::create(builder, loc, builder.getI32IntegerAttr(0), {}); } @@ -9517,7 +9517,7 @@ static void genTMABulkStore(fir::FirOpBuilder &builder, mlir::Location loc, mlir::NVVM::CpAsyncBulkSharedCTAToGlobalOp::create(builder, loc, dst, src, size, {}, {}); mlir::NVVM::InlinePtxOp::create(builder, loc, mlir::TypeRange{}, {}, {}, - "cp.async.bulk.commit_group", {}); + "cp.async.bulk.commit_group;", {}); mlir::NVVM::CpAsyncBulkWaitGroupOp::create(builder, loc, builder.getI32IntegerAttr(0), {}); } diff --git a/flang/test/Lower/CUDA/cuda-device-proc.cuf b/flang/test/Lower/CUDA/cuda-device-proc.cuf index ed015df..666c394 100644 --- a/flang/test/Lower/CUDA/cuda-device-proc.cuf +++ b/flang/test/Lower/CUDA/cuda-device-proc.cuf @@ -494,7 +494,7 @@ end subroutine ! CHECK-LABEL: func.func @_QPtest_bulk_s2g ! CHECL: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> -! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;" ! CHECK: nvvm.cp.async.bulk.wait_group 0 attributes(device) subroutine testAtomicCasLoop(aa, n) @@ -675,7 +675,7 @@ end subroutine ! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_c4 ! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> -! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;" ! CHECK: nvvm.cp.async.bulk.wait_group 0 attributes(global) subroutine test_tma_bulk_store_c8(c, n) @@ -688,7 +688,7 @@ end subroutine ! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_c8 ! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> -! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;" ! CHECK: nvvm.cp.async.bulk.wait_group 0 attributes(global) subroutine test_tma_bulk_store_i4(c, n) @@ -701,7 +701,7 @@ end subroutine ! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_i4 ! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> -! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;" ! CHECK: nvvm.cp.async.bulk.wait_group 0 attributes(global) subroutine test_tma_bulk_store_i8(c, n) @@ -714,7 +714,7 @@ end subroutine ! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_i8 ! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> -! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;" ! CHECK: nvvm.cp.async.bulk.wait_group 0 @@ -728,7 +728,7 @@ end subroutine ! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_r2 ! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> -! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;" ! CHECK: nvvm.cp.async.bulk.wait_group 0 attributes(global) subroutine test_tma_bulk_store_r4(c, n) @@ -741,7 +741,7 @@ end subroutine ! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_r4 ! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> -! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;" ! CHECK: nvvm.cp.async.bulk.wait_group 0 attributes(global) subroutine test_tma_bulk_store_r8(c, n) @@ -754,5 +754,5 @@ end subroutine ! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_r8 ! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> -! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;" ! CHECK: nvvm.cp.async.bulk.wait_group 0 |