[flang-commits] [flang] [Flang] Generalize intrinsic CUBLAS USE association (PR #220024)
Andre Kuhlenschmidt via flang-commits
flang-commits at lists.llvm.org
Mon Oct 5 14:56:27 PDT 2026
https://github.com/akuhlens updated https://github.com/llvm/llvm-project/pull/220024
>From 0e72b65e26cc4e9bd7a144f61306603975a8dde9 Mon Sep 17 00:00:00 2001
From: Andre Kuhlenschmidt <akuhlenschmi at nvidia.com>
Date: Fri, 28 Aug 2026 16:23:30 -0700
Subject: [PATCH 1/4] [Flang] Generalize intrinsic CUBLAS USE association
---
flang/lib/Semantics/resolve-names.cpp | 25 +++---
.../CUDA/cuf-use-cublas-blas-generic.cuf | 86 +++++++++++++++++++
2 files changed, 100 insertions(+), 11 deletions(-)
create mode 100644 flang/test/Semantics/CUDA/cuf-use-cublas-blas-generic.cuf
diff --git a/flang/lib/Semantics/resolve-names.cpp b/flang/lib/Semantics/resolve-names.cpp
index 67690d6f47d06..74b33301b1bb7 100644
--- a/flang/lib/Semantics/resolve-names.cpp
+++ b/flang/lib/Semantics/resolve-names.cpp
@@ -4352,23 +4352,29 @@ static bool HasCUDADummyDataAttribute(const Symbol &procedure) {
struct IntrinsicModuleUseAssociationRule {
const char *moduleName;
- const char *genericName;
- bool (*matches)(SemanticsContext &, const GenericDetails &, const Symbol &);
+ bool (*matches)(SemanticsContext &, const Symbol &, const Symbol &);
};
-static bool MatchesCublasGemm(SemanticsContext &context,
- const GenericDetails &generic, const Symbol &other) {
- const Symbol *specific{generic.specific()};
+static bool MatchesCublasBlas(
+ SemanticsContext &context, const Symbol &generic, const Symbol &other) {
+ if (generic.GetUltimate().name().ToString().rfind("cublas", 0) == 0) {
+ return false;
+ }
+ const auto &details{generic.get<GenericDetails>()};
+ const Symbol *specific{details.specific()};
if (!specific ||
!AreSameProcedureForUseAssociation(context, *specific, other)) {
return false;
}
bool containsSpecific{false};
bool hasCUDAOverload{false};
- for (const Symbol &candidate : generic.specificProcs()) {
+ for (const Symbol &candidate : details.specificProcs()) {
containsSpecific |= &candidate.GetUltimate() == &specific->GetUltimate();
hasCUDAOverload |= HasCUDADummyDataAttribute(candidate);
}
+ // If a CUBLAS generic is found whose CUDA specifics are not compatible
+ // enough with its homonymous host specific, also consider checking their
+ // basic procedure shapes here while ignoring CUDA-specific attributes.
return containsSpecific && hasCUDAOverload;
}
@@ -4378,9 +4384,7 @@ FindIntrinsicModuleUseAssociationRule(
// Add entries here for intrinsic module generics that should take precedence
// over an equivalent external interface during USE association.
static const IntrinsicModuleUseAssociationRule rules[]{
- {"cublas", "sgemm", MatchesCublasGemm},
- {"cublas", "dgemm", MatchesCublasGemm},
- {"cublas", "zgemm", MatchesCublasGemm},
+ {"cublas", MatchesCublasBlas},
};
const Scope &owner{generic.GetUltimate().owner()};
if (!owner.IsModule() || !owner.parent().IsIntrinsicModules() ||
@@ -4389,8 +4393,7 @@ FindIntrinsicModuleUseAssociationRule(
}
for (const auto &rule : rules) {
if (owner.GetName().value() == rule.moduleName &&
- generic.GetUltimate().name() == rule.genericName &&
- rule.matches(context, generic.get<GenericDetails>(), other)) {
+ rule.matches(context, generic, other)) {
return &rule;
}
}
diff --git a/flang/test/Semantics/CUDA/cuf-use-cublas-blas-generic.cuf b/flang/test/Semantics/CUDA/cuf-use-cublas-blas-generic.cuf
new file mode 100644
index 0000000000000..fba20678a7383
--- /dev/null
+++ b/flang/test/Semantics/CUDA/cuf-use-cublas-blas-generic.cuf
@@ -0,0 +1,86 @@
+! RUN: split-file %s %t
+! RUN: %flang_fc1 -emit-obj -x cuda -module-dir %t %t/cublas.cuf -o %t/cublas.o
+! RUN: %flang_fc1 -emit-obj -module-dir %t %t/blas.f90 -o %t/blas.o
+! RUN: %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/external-first.cuf 2>&1 | FileCheck --check-prefix=UNPREFIXED %s
+! RUN: %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/cublas-first.cuf 2>&1 | FileCheck --check-prefix=UNPREFIXED %s
+! RUN: not %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/cublas-prefixed.cuf 2>&1 | FileCheck --check-prefix=CUBLAS-PREFIXED %s
+
+! Unprefixed generics in the intrinsic CUBLAS module model the legacy BLAS
+! interface. Compatible external interfaces may not override those generics.
+! CUBLAS-prefixed API generics do not receive this compatibility treatment.
+
+!--- cublas.cuf
+module cublas
+ implicit none
+ interface
+ subroutine saxpy(n, x)
+ integer :: n
+ real :: x(*)
+ end subroutine
+ subroutine saxpycu(n, x)
+ integer, value :: n
+ real, device :: x(*)
+ end subroutine
+ subroutine cublasaxpy(n, x)
+ integer :: n
+ real :: x(*)
+ end subroutine
+ subroutine cublasaxpycu(n, x)
+ integer, value :: n
+ real, device :: x(*)
+ end subroutine
+ end interface
+ interface saxpy
+ procedure :: saxpy, saxpycu
+ end interface
+ interface cublasaxpy
+ procedure :: cublasaxpy, cublasaxpycu
+ end interface
+end module
+
+!--- blas.f90
+module blas
+ implicit none
+ interface
+ subroutine saxpy(n, x)
+ integer :: n
+ real :: x(*)
+ end subroutine
+ subroutine cublasaxpy(n, x)
+ integer :: n
+ real :: x(*)
+ end subroutine
+ end interface
+end module
+
+!--- external-first.cuf
+module external_first
+ use blas, only: saxpy
+ use, intrinsic :: cublas, only: saxpy
+contains
+ subroutine test_cuda_specific()
+ real, device :: x(1)
+ call saxpy(1, x)
+ end subroutine
+end module
+
+!--- cublas-first.cuf
+module cublas_first
+ use, intrinsic :: cublas, only: saxpy
+ use blas, only: saxpy
+contains
+ subroutine test_cuda_specific()
+ real, device :: x(1)
+ call saxpy(1, x)
+ end subroutine
+end module
+
+!--- cublas-prefixed.cuf
+module cublas_prefixed
+ use blas, only: cublasaxpy
+ use, intrinsic :: cublas, only: cublasaxpy
+ procedure(cublasaxpy), pointer :: p
+end module
+
+! UNPREFIXED: warning: USE association selects intrinsic 'cublas' generic 'saxpy' over an equivalent external interface [-Wprefer-intrinsic-module-use-association]
+! CUBLAS-PREFIXED: error: 'cublasaxpy' must be an abstract interface or a procedure with an explicit interface
>From d266d754d86bace27fcd1373cfc48cc586266f17 Mon Sep 17 00:00:00 2001
From: Andre Kuhlenschmidt <akuhlenschmi at nvidia.com>
Date: Tue, 1 Sep 2026 16:07:46 -0700
Subject: [PATCH 2/4] [Flang] Test disjoint CUBLAS USE ONLY imports
---
.../CUDA/cuf-use-cublas-blas-generic.cuf | 30 +++++++++++++++++++
1 file changed, 30 insertions(+)
diff --git a/flang/test/Semantics/CUDA/cuf-use-cublas-blas-generic.cuf b/flang/test/Semantics/CUDA/cuf-use-cublas-blas-generic.cuf
index fba20678a7383..901e0a77078ac 100644
--- a/flang/test/Semantics/CUDA/cuf-use-cublas-blas-generic.cuf
+++ b/flang/test/Semantics/CUDA/cuf-use-cublas-blas-generic.cuf
@@ -4,6 +4,8 @@
! RUN: %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/external-first.cuf 2>&1 | FileCheck --check-prefix=UNPREFIXED %s
! RUN: %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/cublas-first.cuf 2>&1 | FileCheck --check-prefix=UNPREFIXED %s
! RUN: not %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/cublas-prefixed.cuf 2>&1 | FileCheck --check-prefix=CUBLAS-PREFIXED %s
+! RUN: %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/blas-first-disjoint-only.cuf 2>&1 | FileCheck --allow-empty --check-prefix=DISJOINT-ONLY %s
+! RUN: %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/cublas-first-disjoint-only.cuf 2>&1 | FileCheck --allow-empty --check-prefix=DISJOINT-ONLY %s
! Unprefixed generics in the intrinsic CUBLAS module model the legacy BLAS
! interface. Compatible external interfaces may not override those generics.
@@ -82,5 +84,33 @@ module cublas_prefixed
procedure(cublasaxpy), pointer :: p
end module
+! The ONLY lists below are disjoint even though both used modules define both
+! names, so no intrinsic-module preference rule should be needed.
+
+!--- blas-first-disjoint-only.cuf
+module blas_first_disjoint_only
+ use blas, only: saxpy
+ use, intrinsic :: cublas, only: cublasaxpy
+ procedure(saxpy), pointer :: p
+contains
+ subroutine test_cuda_specific()
+ real, device :: x(1)
+ call cublasaxpy(1, x)
+ end subroutine
+end module
+
+!--- cublas-first-disjoint-only.cuf
+module cublas_first_disjoint_only
+ use, intrinsic :: cublas, only: cublasaxpy
+ use blas, only: saxpy
+ procedure(saxpy), pointer :: p
+contains
+ subroutine test_cuda_specific()
+ real, device :: x(1)
+ call cublasaxpy(1, x)
+ end subroutine
+end module
+
! UNPREFIXED: warning: USE association selects intrinsic 'cublas' generic 'saxpy' over an equivalent external interface [-Wprefer-intrinsic-module-use-association]
! CUBLAS-PREFIXED: error: 'cublasaxpy' must be an abstract interface or a procedure with an explicit interface
+! DISJOINT-ONLY-NOT: {{warning:|error:}}
>From b2d929fa655dd6f02e04bec361c3a8a17e785c4c Mon Sep 17 00:00:00 2001
From: Andre Kuhlenschmidt <akuhlenschmi at nvidia.com>
Date: Mon, 5 Oct 2026 10:39:59 -0700
Subject: [PATCH 3/4] [Flang] Address CUBLAS USE association review
---
flang/docs/Extensions.md | 14 +
flang/lib/Semantics/resolve-names.cpp | 17 +-
.../CUDA/cuf-use-cublas-blas-generic.cuf | 244 +++++++++++++++++-
3 files changed, 268 insertions(+), 7 deletions(-)
diff --git a/flang/docs/Extensions.md b/flang/docs/Extensions.md
index 64f502ee6e484..a43ede23f816e 100644
--- a/flang/docs/Extensions.md
+++ b/flang/docs/Extensions.md
@@ -461,6 +461,19 @@ print *, is_contiguous(a(::2)) ! prints T in Flang
* When a name is brought into a scope by multiple ways,
such as USE-association as well as an `IMPORT` from its host,
it's an error only if the resolution is ambiguous.
+* When USE association brings an equivalent external procedure interface
+ and a legacy BLAS generic from the intrinsic `cublas` or `cublas_v2`
+ module into the same scope under the same local name, Flang selects
+ the intrinsic module's generic, regardless of USE statement order.
+ The generic must contain a same-named host specific with characteristics
+ equal to those of the external interface, and at least one specific
+ with a CUDA dummy data attribute.
+ Generic names beginning with `cublas` are excluded.
+ A warning is emitted by default, including with `-pedantic`.
+ Use `-Wno-prefer-intrinsic-module-use-association` to suppress the warning,
+ or `-fno-prefer-intrinsic-module-use-association` to disable the extension.
+ This extension is deprecated and may be removed at any time; it was added
+ to support BerkeleyGW.
* An entity may appear in a `DATA` statement before its explicit
type declaration under `IMPLICIT NONE(TYPE)`.
* `INCLUDE` lines can start in any column, can be preceded in
@@ -650,6 +663,7 @@ end program
multiple modules, the name must refer to a generic interface; PGI
allows a name to be a procedure from one module and a generic interface
from another.
+ Flang supports the limited intrinsic CUBLAS exception described above.
* Type parameter declarations must come first in a derived type definition;
some compilers allow them to follow `PRIVATE`, or be intermixed with the
component declarations.
diff --git a/flang/lib/Semantics/resolve-names.cpp b/flang/lib/Semantics/resolve-names.cpp
index 74b33301b1bb7..d1ad3e905942d 100644
--- a/flang/lib/Semantics/resolve-names.cpp
+++ b/flang/lib/Semantics/resolve-names.cpp
@@ -4301,6 +4301,14 @@ static bool CheckCompatibleDistinctUltimates(SemanticsContext &context,
return true; // don't try to merge generics (or whatever)
}
+// Check whether two symbols identify the same procedure. Otherwise, require
+// matching names and procedure kinds: intrinsic procedures must both be
+// intrinsic, module procedures must identify the same module symbol, and
+// external procedures must have explicit interfaces with equal characteristics.
+// For intrinsic-module compatibility rules, a future, separate check could
+// compare the basic shapes of host and CUDA specifics, including dummy argument
+// and result types and ranks, while ignoring CUDA-specific attributes.
+// That shape comparison is not performed here.
static bool AreSameProcedureForUseAssociation(
SemanticsContext &context, const Symbol &p1, const Symbol &p2) {
const Symbol &ultimate1{p1.GetUltimate()};
@@ -4357,7 +4365,10 @@ struct IntrinsicModuleUseAssociationRule {
static bool MatchesCublasBlas(
SemanticsContext &context, const Symbol &generic, const Symbol &other) {
- if (generic.GetUltimate().name().ToString().rfind("cublas", 0) == 0) {
+ // Exclude CUBLAS-prefixed API names from the legacy BLAS compatibility rule.
+ const SourceName &genericName{generic.GetUltimate().name()};
+ if (llvm::StringRef{genericName.begin(), genericName.size()}.starts_with(
+ "cublas")) {
return false;
}
const auto &details{generic.get<GenericDetails>()};
@@ -4372,9 +4383,6 @@ static bool MatchesCublasBlas(
containsSpecific |= &candidate.GetUltimate() == &specific->GetUltimate();
hasCUDAOverload |= HasCUDADummyDataAttribute(candidate);
}
- // If a CUBLAS generic is found whose CUDA specifics are not compatible
- // enough with its homonymous host specific, also consider checking their
- // basic procedure shapes here while ignoring CUDA-specific attributes.
return containsSpecific && hasCUDAOverload;
}
@@ -4385,6 +4393,7 @@ FindIntrinsicModuleUseAssociationRule(
// over an equivalent external interface during USE association.
static const IntrinsicModuleUseAssociationRule rules[]{
{"cublas", MatchesCublasBlas},
+ {"cublas_v2", MatchesCublasBlas},
};
const Scope &owner{generic.GetUltimate().owner()};
if (!owner.IsModule() || !owner.parent().IsIntrinsicModules() ||
diff --git a/flang/test/Semantics/CUDA/cuf-use-cublas-blas-generic.cuf b/flang/test/Semantics/CUDA/cuf-use-cublas-blas-generic.cuf
index 901e0a77078ac..c13576295fbe9 100644
--- a/flang/test/Semantics/CUDA/cuf-use-cublas-blas-generic.cuf
+++ b/flang/test/Semantics/CUDA/cuf-use-cublas-blas-generic.cuf
@@ -1,13 +1,36 @@
! RUN: split-file %s %t
-! RUN: %flang_fc1 -emit-obj -x cuda -module-dir %t %t/cublas.cuf -o %t/cublas.o
-! RUN: %flang_fc1 -emit-obj -module-dir %t %t/blas.f90 -o %t/blas.o
+! RUN: %flang_fc1 -fsyntax-only -x cuda -module-dir %t %t/cublas.cuf
+! RUN: %flang_fc1 -fsyntax-only -module-dir %t %t/blas.f90
+! RUN: %flang_fc1 -fsyntax-only -x cuda -module-dir %t %t/cublas_v2.cuf
+! RUN: %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/v2-external-first.cuf 2>&1 | FileCheck --check-prefix=V2-UNPREFIXED %s
+! RUN: %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/v2-cublas-first.cuf 2>&1 | FileCheck --check-prefix=V2-UNPREFIXED %s
+! RUN: not %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/v2-cublas-prefixed.cuf 2>&1 | FileCheck --check-prefix=CUBLAS-PREFIXED %s
! RUN: %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/external-first.cuf 2>&1 | FileCheck --check-prefix=UNPREFIXED %s
! RUN: %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/cublas-first.cuf 2>&1 | FileCheck --check-prefix=UNPREFIXED %s
! RUN: not %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/cublas-prefixed.cuf 2>&1 | FileCheck --check-prefix=CUBLAS-PREFIXED %s
! RUN: %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/blas-first-disjoint-only.cuf 2>&1 | FileCheck --allow-empty --check-prefix=DISJOINT-ONLY %s
! RUN: %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/cublas-first-disjoint-only.cuf 2>&1 | FileCheck --allow-empty --check-prefix=DISJOINT-ONLY %s
-! Unprefixed generics in the intrinsic CUBLAS module model the legacy BLAS
+! RUN: %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/function-external-first.cuf 2>&1 | FileCheck --check-prefix=FUNCTION %s
+! RUN: %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/function-cublas-first.cuf 2>&1 | FileCheck --check-prefix=FUNCTION %s
+! RUN: %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/v2-function-external-first.cuf 2>&1 | FileCheck --check-prefix=V2-FUNCTION %s
+! RUN: %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/v2-function-cublas-first.cuf 2>&1 | FileCheck --check-prefix=V2-FUNCTION %s
+
+! RUN: not %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/no-cuda-external-first.cuf 2>&1 | FileCheck --check-prefix=NO-CUDA %s
+! RUN: not %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/no-cuda-cublas-first.cuf 2>&1 | FileCheck --check-prefix=NO-CUDA %s
+! RUN: not %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/v2-no-cuda-external-first.cuf 2>&1 | FileCheck --check-prefix=NO-CUDA %s
+! RUN: not %flang_fc1 -fsyntax-only -x cuda -I %t -fintrinsic-modules-path %t %t/v2-no-cuda-cublas-first.cuf 2>&1 | FileCheck --check-prefix=NO-CUDA %s
+
+! RUN: %flang_fc1 -emit-hlfir -x cuda -I %t -fintrinsic-modules-path %t %t/external-first.cuf -o - | FileCheck --check-prefix=HLFIR-SUBROUTINE %s
+! RUN: %flang_fc1 -emit-hlfir -x cuda -I %t -fintrinsic-modules-path %t %t/cublas-first.cuf -o - | FileCheck --check-prefix=HLFIR-SUBROUTINE %s
+! RUN: %flang_fc1 -emit-hlfir -x cuda -I %t -fintrinsic-modules-path %t %t/function-external-first.cuf -o - | FileCheck --check-prefix=HLFIR-FUNCTION %s
+! RUN: %flang_fc1 -emit-hlfir -x cuda -I %t -fintrinsic-modules-path %t %t/function-cublas-first.cuf -o - | FileCheck --check-prefix=HLFIR-FUNCTION %s
+! RUN: %flang_fc1 -emit-hlfir -x cuda -I %t -fintrinsic-modules-path %t %t/v2-external-first.cuf -o - | FileCheck --check-prefix=HLFIR-SUBROUTINE %s
+! RUN: %flang_fc1 -emit-hlfir -x cuda -I %t -fintrinsic-modules-path %t %t/v2-cublas-first.cuf -o - | FileCheck --check-prefix=HLFIR-SUBROUTINE %s
+! RUN: %flang_fc1 -emit-hlfir -x cuda -I %t -fintrinsic-modules-path %t %t/v2-function-external-first.cuf -o - | FileCheck --check-prefix=HLFIR-FUNCTION %s
+! RUN: %flang_fc1 -emit-hlfir -x cuda -I %t -fintrinsic-modules-path %t %t/v2-function-cublas-first.cuf -o - | FileCheck --check-prefix=HLFIR-FUNCTION %s
+
+! Unprefixed generics in the intrinsic CUBLAS modules model the legacy BLAS
! interface. Compatible external interfaces may not override those generics.
! CUBLAS-prefixed API generics do not receive this compatibility treatment.
@@ -19,6 +42,10 @@ module cublas
integer :: n
real :: x(*)
end subroutine
+ subroutine scopy(n, x)
+ integer :: n
+ real :: x(*)
+ end subroutine
subroutine saxpycu(n, x)
integer, value :: n
real, device :: x(*)
@@ -31,23 +58,94 @@ module cublas
integer, value :: n
real, device :: x(*)
end subroutine
+ integer function isamax(n, x)
+ integer :: n
+ real :: x(*)
+ end function
+ integer function isamaxcu(n, x)
+ integer, value :: n
+ real, device :: x(*)
+ end function
end interface
interface saxpy
procedure :: saxpy, saxpycu
end interface
+ ! This host-only generic must not receive compatibility treatment.
+ interface scopy
+ procedure :: scopy
+ end interface
interface cublasaxpy
procedure :: cublasaxpy, cublasaxpycu
end interface
+ interface isamax
+ procedure :: isamax, isamaxcu
+ end interface
+end module
+
+!--- cublas_v2.cuf
+module cublas_v2
+ implicit none
+ interface
+ subroutine saxpy(n, x)
+ integer :: n
+ real :: x(*)
+ end subroutine
+ subroutine scopy(n, x)
+ integer :: n
+ real :: x(*)
+ end subroutine
+ subroutine saxpycu(n, x)
+ integer, value :: n
+ real, device :: x(*)
+ end subroutine
+ subroutine cublasaxpy(n, x)
+ integer :: n
+ real :: x(*)
+ end subroutine
+ subroutine cublasaxpycu(n, x)
+ integer, value :: n
+ real, device :: x(*)
+ end subroutine
+ integer function isamax(n, x)
+ integer :: n
+ real :: x(*)
+ end function
+ integer function isamaxcu(n, x)
+ integer, value :: n
+ real, device :: x(*)
+ end function
+ end interface
+ interface saxpy
+ procedure :: saxpy, saxpycu
+ end interface
+ ! This host-only generic must not receive compatibility treatment.
+ interface scopy
+ procedure :: scopy
+ end interface
+ interface cublasaxpy
+ procedure :: cublasaxpy, cublasaxpycu
+ end interface
+ interface isamax
+ procedure :: isamax, isamaxcu
+ end interface
end module
!--- blas.f90
module blas
implicit none
interface
+ integer function isamax(n, x)
+ integer :: n
+ real :: x(*)
+ end function
subroutine saxpy(n, x)
integer :: n
real :: x(*)
end subroutine
+ subroutine scopy(n, x)
+ integer :: n
+ real :: x(*)
+ end subroutine
subroutine cublasaxpy(n, x)
integer :: n
real :: x(*)
@@ -84,6 +182,127 @@ module cublas_prefixed
procedure(cublasaxpy), pointer :: p
end module
+!--- v2-external-first.cuf
+module v2_external_first
+ use blas, only: saxpy
+ use, intrinsic :: cublas_v2, only: saxpy
+contains
+ subroutine test_cuda_specific()
+ real, device :: x(1)
+ call saxpy(1, x)
+ end subroutine
+end module
+
+!--- v2-cublas-first.cuf
+module v2_cublas_first
+ use, intrinsic :: cublas_v2, only: saxpy
+ use blas, only: saxpy
+contains
+ subroutine test_cuda_specific()
+ real, device :: x(1)
+ call saxpy(1, x)
+ end subroutine
+end module
+
+!--- v2-cublas-prefixed.cuf
+module v2_cublas_prefixed
+ use blas, only: cublasaxpy
+ use, intrinsic :: cublas_v2, only: cublasaxpy
+ procedure(cublasaxpy), pointer :: p
+end module
+
+!--- function-external-first.cuf
+module function_external_first
+ use blas, only: isamax
+ use, intrinsic :: cublas, only: isamax
+contains
+ subroutine test_cuda_function()
+ integer :: k
+ real, device :: x(4)
+ k = isamax(4, x)
+ end subroutine
+end module
+
+!--- function-cublas-first.cuf
+module function_cublas_first
+ use, intrinsic :: cublas, only: isamax
+ use blas, only: isamax
+contains
+ subroutine test_cuda_function()
+ integer :: k
+ real, device :: x(4)
+ k = isamax(4, x)
+ end subroutine
+end module
+
+!--- v2-function-external-first.cuf
+module v2_function_external_first
+ use blas, only: isamax
+ use, intrinsic :: cublas_v2, only: isamax
+contains
+ subroutine test_cuda_function()
+ integer :: k
+ real, device :: x(4)
+ k = isamax(4, x)
+ end subroutine
+end module
+
+!--- v2-function-cublas-first.cuf
+module v2_function_cublas_first
+ use, intrinsic :: cublas_v2, only: isamax
+ use blas, only: isamax
+contains
+ subroutine test_cuda_function()
+ integer :: k
+ real, device :: x(4)
+ k = isamax(4, x)
+ end subroutine
+end module
+
+!--- no-cuda-external-first.cuf
+module no_cuda_external_first
+ use blas, only: scopy
+ use, intrinsic :: cublas, only: scopy
+contains
+ subroutine test_host_only_generic()
+ real :: x(4)
+ call scopy(4, x)
+ end subroutine
+end module
+
+!--- no-cuda-cublas-first.cuf
+module no_cuda_cublas_first
+ use, intrinsic :: cublas, only: scopy
+ use blas, only: scopy
+contains
+ subroutine test_host_only_generic()
+ real :: x(4)
+ call scopy(4, x)
+ end subroutine
+end module
+
+!--- v2-no-cuda-external-first.cuf
+module v2_no_cuda_external_first
+ use blas, only: scopy
+ use, intrinsic :: cublas_v2, only: scopy
+contains
+ subroutine test_host_only_generic()
+ real :: x(4)
+ call scopy(4, x)
+ end subroutine
+end module
+
+!--- v2-no-cuda-cublas-first.cuf
+module v2_no_cuda_cublas_first
+ use, intrinsic :: cublas_v2, only: scopy
+ use blas, only: scopy
+contains
+ subroutine test_host_only_generic()
+ real :: x(4)
+ call scopy(4, x)
+ end subroutine
+end module
+
! The ONLY lists below are disjoint even though both used modules define both
! names, so no intrinsic-module preference rule should be needed.
@@ -112,5 +331,24 @@ contains
end module
! UNPREFIXED: warning: USE association selects intrinsic 'cublas' generic 'saxpy' over an equivalent external interface [-Wprefer-intrinsic-module-use-association]
+! V2-UNPREFIXED: warning: USE association selects intrinsic 'cublas_v2' generic 'saxpy' over an equivalent external interface [-Wprefer-intrinsic-module-use-association]
! CUBLAS-PREFIXED: error: 'cublasaxpy' must be an abstract interface or a procedure with an explicit interface
! DISJOINT-ONLY-NOT: {{warning:|error:}}
+! FUNCTION: warning: USE association selects intrinsic 'cublas' generic 'isamax' over an equivalent external interface [-Wprefer-intrinsic-module-use-association]
+! V2-FUNCTION: warning: USE association selects intrinsic 'cublas_v2' generic 'isamax' over an equivalent external interface [-Wprefer-intrinsic-module-use-association]
+! NO-CUDA-NOT: warning: USE association selects intrinsic
+! NO-CUDA: error: Reference to 'scopy' is ambiguous
+! NO-CUDA-NOT: warning: USE association selects intrinsic
+
+! Device calls must bind to the CUDA specifics of the selected intrinsic generic.
+! HLFIR-SUBROUTINE-LABEL: func.func @_QM{{.*}}Ptest_cuda_specific()
+! HLFIR-SUBROUTINE-NOT: fir.call @_QPsaxpy(
+! HLFIR-SUBROUTINE: fir.call @_QPsaxpycu(
+! HLFIR-SUBROUTINE-NOT: fir.call @_QPsaxpy(
+! HLFIR-SUBROUTINE: return
+! HLFIR-FUNCTION-LABEL: func.func @_QM{{.*}}Ptest_cuda_function()
+! HLFIR-FUNCTION-NOT: fir.call @_QPisamax(
+! HLFIR-FUNCTION: %[[RESULT:.*]] = fir.call @_QPisamaxcu({{.*}}) {{.*}} -> i32
+! HLFIR-FUNCTION: hlfir.assign %[[RESULT]] to
+! HLFIR-FUNCTION-NOT: fir.call @_QPisamax(
+! HLFIR-FUNCTION: return
>From f1b785a00deb7a15d5fb79894719d4ac1e611958 Mon Sep 17 00:00:00 2001
From: Andre Kuhlenschmidt <akuhlenschmi at nvidia.com>
Date: Mon, 5 Oct 2026 14:51:48 -0700
Subject: [PATCH 4/4] [Flang][NFC] Name CUBLAS generic string view
---
flang/lib/Semantics/resolve-names.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/flang/lib/Semantics/resolve-names.cpp b/flang/lib/Semantics/resolve-names.cpp
index d1ad3e905942d..cde4cafc6dba2 100644
--- a/flang/lib/Semantics/resolve-names.cpp
+++ b/flang/lib/Semantics/resolve-names.cpp
@@ -4367,8 +4367,8 @@ static bool MatchesCublasBlas(
SemanticsContext &context, const Symbol &generic, const Symbol &other) {
// Exclude CUBLAS-prefixed API names from the legacy BLAS compatibility rule.
const SourceName &genericName{generic.GetUltimate().name()};
- if (llvm::StringRef{genericName.begin(), genericName.size()}.starts_with(
- "cublas")) {
+ const llvm::StringRef genericNameRef{genericName.begin(), genericName.size()};
+ if (genericNameRef.starts_with("cublas")) {
return false;
}
const auto &details{generic.get<GenericDetails>()};
More information about the flang-commits
mailing list