[flang-commits] [flang] [flang][cuda] Add bind c to cudadevice procedures (PR #92822)
Valentin Clement バレンタイン クレメン via flang-commits
flang-commits at lists.llvm.org
Tue May 28 11:54:22 PDT 2024
https://github.com/clementval updated https://github.com/llvm/llvm-project/pull/92822
>From f044e73f6fc0a46e74ae24d761f4277bc6c6ab54 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Mon, 20 May 2024 14:17:45 -0700
Subject: [PATCH 1/3] [flang][cuda] Add bind c to cudadevice procedures
---
flang/module/cudadevice.f90 | 16 +++++-----
flang/test/Lower/CUDA/cuda-device-proc.cuf | 36 ++++++++++++++++++++++
2 files changed, 44 insertions(+), 8 deletions(-)
create mode 100644 flang/test/Lower/CUDA/cuda-device-proc.cuf
diff --git a/flang/module/cudadevice.f90 b/flang/module/cudadevice.f90
index f34820dd10792..5770701cf24d4 100644
--- a/flang/module/cudadevice.f90
+++ b/flang/module/cudadevice.f90
@@ -18,34 +18,34 @@ module cudadevice
! Synchronization Functions
interface
- attributes(device) subroutine syncthreads()
+ attributes(device) subroutine syncthreads() bind(c, name='__syncthreads')
end subroutine
end interface
public :: syncthreads
interface
- attributes(device) integer function syncthreads_and(value)
+ attributes(device) integer function syncthreads_and(value) bind(c, name='__syncthreads_and')
integer :: value
end function
end interface
public :: syncthreads_and
interface
- attributes(device) integer function syncthreads_count(value)
+ attributes(device) integer function syncthreads_count(value) bind(c, name='__syncthreads_count')
integer :: value
end function
end interface
public :: syncthreads_count
interface
- attributes(device) integer function syncthreads_or(value)
+ attributes(device) integer function syncthreads_or(value) bind(c, name='__syncthreads_or')
integer :: value
end function
end interface
public :: syncthreads_or
interface
- attributes(device) subroutine syncwarp(mask)
+ attributes(device) subroutine syncwarp(mask) bind(c, name='__syncwrap')
integer :: mask
end subroutine
end interface
@@ -54,19 +54,19 @@ attributes(device) subroutine syncwarp(mask)
! Memory Fences
interface
- attributes(device) subroutine threadfence()
+ attributes(device) subroutine threadfence() bind(c, name='__threadfence')
end subroutine
end interface
public :: threadfence
interface
- attributes(device) subroutine threadfence_block()
+ attributes(device) subroutine threadfence_block() bind(c, name='__threadfence_block')
end subroutine
end interface
public :: threadfence_block
interface
- attributes(device) subroutine threadfence_system()
+ attributes(device) subroutine threadfence_system() bind(c, name='__threadfence_system')
end subroutine
end interface
public :: threadfence_system
diff --git a/flang/test/Lower/CUDA/cuda-device-proc.cuf b/flang/test/Lower/CUDA/cuda-device-proc.cuf
new file mode 100644
index 0000000000000..e890f81d9a238
--- /dev/null
+++ b/flang/test/Lower/CUDA/cuda-device-proc.cuf
@@ -0,0 +1,36 @@
+! RUN: bbc -emit-hlfir -fcuda %s -o - | FileCheck %s
+
+! Test CUDA Fortran procedures available in cudadevice module
+
+attributes(global) subroutine devsub()
+ implicit none
+ integer :: ret
+
+ call syncthreads()
+ call syncwarp(1)
+ call threadfence()
+ call threadfence_block()
+ call threadfence_system()
+ ret = syncthreads_and(1)
+ ret = syncthreads_count(1)
+ ret = syncthreads_or(1)
+end
+
+! CHECK-LABEL: func.func @_QPdevsub() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}
+! CHECK: fir.call @__syncthreads()
+! CHECK: fir.call @__syncwrap(%{{.*}}) fastmath<contract> : (!fir.ref<i32>) -> ()
+! CHECK: fir.call @__threadfence()
+! CHECK: fir.call @__threadfence_block()
+! CHECK: fir.call @__threadfence_system()
+! CHECK: %{{.*}} = fir.call @__syncthreads_and(%{{.*}}) fastmath<contract> : (!fir.ref<i32>) -> i32
+! CHECK: %{{.*}} = fir.call @__syncthreads_count(%{{.*}}) fastmath<contract> : (!fir.ref<i32>) -> i32
+! CHECK: %{{.*}} = fir.call @__syncthreads_or(%{{.*}}) fastmath<contract> : (!fir.ref<i32>) -> i32
+
+! CHECK: func.func private @__syncthreads() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncthreads"}
+! CHECK: func.func private @__syncwrap(!fir.ref<i32> {cuf.data_attr = #cuf.cuda<device>}) attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncwrap"}
+! CHECK: func.func private @__threadfence() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__threadfence"}
+! CHECK: func.func private @__threadfence_block() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__threadfence_block"}
+! CHECK: func.func private @__threadfence_system() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__threadfence_system"}
+! CHECK: func.func private @__syncthreads_and(!fir.ref<i32> {cuf.data_attr = #cuf.cuda<device>}) -> i32 attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncthreads_and"}
+! CHECK: func.func private @__syncthreads_count(!fir.ref<i32> {cuf.data_attr = #cuf.cuda<device>}) -> i32 attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncthreads_count"}
+! CHECK: func.func private @__syncthreads_or(!fir.ref<i32> {cuf.data_attr = #cuf.cuda<device>}) -> i32 attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncthreads_or"}
>From fbdf62a9b2ee8ce874949438ce822f8a1a54e198 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Tue, 28 May 2024 08:34:56 -0700
Subject: [PATCH 2/3] Fx typo
---
flang/module/cudadevice.f90 | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/flang/module/cudadevice.f90 b/flang/module/cudadevice.f90
index 5770701cf24d4..0224ecfdde7c6 100644
--- a/flang/module/cudadevice.f90
+++ b/flang/module/cudadevice.f90
@@ -45,7 +45,7 @@ attributes(device) integer function syncthreads_or(value) bind(c, name='__syncth
public :: syncthreads_or
interface
- attributes(device) subroutine syncwarp(mask) bind(c, name='__syncwrap')
+ attributes(device) subroutine syncwarp(mask) bind(c, name='__syncwarp')
integer :: mask
end subroutine
end interface
>From a575399f06b6a8fb6d54ca72f8e7d21c3ce1f5e4 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Valentin=20Clement=20=28=E3=83=90=E3=83=AC=E3=83=B3?=
=?UTF-8?q?=E3=82=BF=E3=82=A4=E3=83=B3=20=E3=82=AF=E3=83=AC=E3=83=A1?=
=?UTF-8?q?=E3=83=B3=29?= <clementval at gmail.com>
Date: Tue, 28 May 2024 11:54:14 -0700
Subject: [PATCH 3/3] Update flang/test/Lower/CUDA/cuda-device-proc.cuf
---
flang/test/Lower/CUDA/cuda-device-proc.cuf | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/flang/test/Lower/CUDA/cuda-device-proc.cuf b/flang/test/Lower/CUDA/cuda-device-proc.cuf
index e890f81d9a238..34297bc897709 100644
--- a/flang/test/Lower/CUDA/cuda-device-proc.cuf
+++ b/flang/test/Lower/CUDA/cuda-device-proc.cuf
@@ -27,7 +27,7 @@ end
! CHECK: %{{.*}} = fir.call @__syncthreads_or(%{{.*}}) fastmath<contract> : (!fir.ref<i32>) -> i32
! CHECK: func.func private @__syncthreads() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncthreads"}
-! CHECK: func.func private @__syncwrap(!fir.ref<i32> {cuf.data_attr = #cuf.cuda<device>}) attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncwrap"}
+! CHECK: func.func private @__syncwarp(!fir.ref<i32> {cuf.data_attr = #cuf.cuda<device>}) attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncwarp"}
! CHECK: func.func private @__threadfence() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__threadfence"}
! CHECK: func.func private @__threadfence_block() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__threadfence_block"}
! CHECK: func.func private @__threadfence_system() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__threadfence_system"}
More information about the flang-commits
mailing list