mirrors/llvm-project
0
0
Fork 0
mirror of https://github.com/llvm/llvm-project.git synced 2025-04-16 19:06:33 +00:00
Code Issues Projects Releases Wiki Activity

[flang][cuda] Add bind c to cudadevice procedures (#92822)

Browse Source 
This patch adds bind c names to functions and subroutines in cudadevice
so they can be lowered and not hit the intrinsic procedure TODOs.
This commit is contained in:
Valentin Clement (バレンタイン クレメン) 2024-05-28 14:57:13 -07:00 committed by GitHub
parent 060b3023e1
commit 00bd2fa198
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
2 changed files with 44 additions and 8 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

16
flang/module/cudadevice.f90
View File

@ -18,34 +18,34 @@ implicit none
! 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='__syncwarp')
integer :: mask
end subroutine
end interface
@ -54,19 +54,19 @@ implicit none
! 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

36
flang/test/Lower/CUDA/cuda-device-proc.cuf Normal file
View File

@ -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 @__syncwarp(%{{.*}}) 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 @__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"}
! 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"}