Skip to content

Commit 63be744

Browse files
Sergei-LebedevSergei-Lebedev
authored
BUILD/CUDA: build device code for supported CUDA arch (#10882)
1 parent 179a0e7 commit 63be744

File tree

1 file changed

+85
-29
lines changed

1 file changed

+85
-29
lines changed

‎config/m4/cuda.m4‎

Lines changed: 85 additions & 29 deletions
Original file line numberDiff line numberDiff line change
@@ -3,6 +3,28 @@
33
# See file LICENSE for terms.
44
#
55

6+
NVCC_CUDA_MIN_REQUIRED_MAJOR=10
7+
NVCC_CUDA_MIN_REQUIRED_MINOR=2
8+
9+
ARCH9_CODE="-gencode=arch=compute_70,code=sm_70"
10+
ARCH10_CODE="-gencode=arch=compute_75,code=sm_75"
11+
ARCH110_CODE="-gencode=arch=compute_80,code=sm_80"
12+
ARCH111_CODE="-gencode=arch=compute_86,code=sm_86"
13+
ARCH120_CODE="-gencode=arch=compute_90,code=sm_90"
14+
ARCH124_CODE="-gencode=arch=compute_89,code=sm_89"
15+
ARCH128_CODE="-gencode=arch=compute_100,code=sm_100 -gencode=arch=compute_120,code=sm_120"
16+
ARCH130_CODE="-gencode=arch=compute_110,code=sm_110"
17+
18+
19+
ARCH9_PTX="-gencode=arch=compute_70,code=compute_70"
20+
ARCH10_PTX=""
21+
ARCH110_PTX="-gencode=arch=compute_80,code=compute_80"
22+
ARCH111_PTX="-gencode=arch=compute_86,code=compute_86"
23+
ARCH120_PTX="-gencode=arch=compute_90,code=compute_90"
24+
ARCH124_PTX="-gencode=arch=compute_90,code=compute_90"
25+
ARCH128_PTX="-gencode=arch=compute_120,code=compute_120"
26+
ARCH130_PTX="-gencode=arch=compute_120,code=compute_120"
27+
628
# Define CUDA language
729
AC_LANG_DEFINE([CUDA], [cuda], [NVCC], [NVCC], [C++], [
830
ac_ext=cu
@@ -13,11 +35,6 @@ AC_LANG_DEFINE([CUDA], [cuda], [NVCC], [NVCC], [C++], [
1335

1436
# Define CUDA language compiler
1537
AC_DEFUN([AC_LANG_COMPILER(CUDA)], [
16-
AC_ARG_WITH([nvcc-gencode],
17-
[AS_HELP_STRING([--with-nvcc-gencode=(OPTS)], [Build for specific GPU architectures])],
18-
[],
19-
[with_nvcc_gencode="-gencode=arch=compute_80,code=sm_80"])
20-
2138
AC_ARG_VAR([NVCC], [nvcc compiler path])
2239
AC_ARG_VAR([NVCCFLAGS], [nvcc compiler flags])
2340
BASE_NVCCFLAGS="$BASE_NVCCFLAGS -g $with_nvcc_gencode"
@@ -31,30 +48,69 @@ AC_DEFUN([AC_LANG_COMPILER(CUDA)], [
3148
# Check for nvcc compiler support
3249
AC_DEFUN([UCX_CUDA_CHECK_NVCC], [
3350
AS_IF([test "x$NVCC" != "x"], [
34-
AC_MSG_CHECKING([$NVCC needs explicit c++11 option])
35-
AC_LANG_PUSH([CUDA])
36-
AC_COMPILE_IFELSE([AC_LANG_SOURCE([[
37-
#if __cplusplus < 201103L
38-
#error missing C++11
39-
#endif
40-
]])],
41-
[AC_MSG_RESULT([no])],
42-
[AC_MSG_RESULT([yes])
43-
BASE_NVCCFLAGS="$BASE_NVCCFLAGS -std=c++11"])
44-
AC_LANG_POP
45-
46-
AC_MSG_CHECKING([$NVCC can compile])
47-
AC_LANG_PUSH([CUDA])
48-
AC_COMPILE_IFELSE([AC_LANG_SOURCE([[
49-
#include <cuda_runtime.h>
50-
__global__ void my_kernel(void) {}
51-
int main(void) { my_kernel<<<1, 1>>>(); return 0; }
52-
]])],
53-
[AC_MSG_RESULT([yes])],
54-
[AC_MSG_RESULT([no])
55-
NVCC=""])
56-
AC_LANG_POP
57-
])
51+
CUDA_MAJOR_VERSION=`$NVCC --version | grep release | sed 's/.*release //' | sed 's/\,.*//' | cut -d "." -f 1`
52+
CUDA_MINOR_VERSION=`$NVCC --version | grep release | sed 's/.*release //' | sed 's/\,.*//' | cut -d "." -f 2`
53+
AC_MSG_RESULT([Detected CUDA version: $CUDA_MAJOR_VERSION.$CUDA_MINOR_VERSION])
54+
AS_IF([test $CUDA_MAJOR_VERSION -lt $NVCC_CUDA_MIN_REQUIRED_MAJOR -o \( $CUDA_MAJOR_VERSION -eq $NVCC_CUDA_MIN_REQUIRED_MAJOR -a $CUDA_MINOR_VERSION -lt $NVCC_CUDA_MIN_REQUIRED_MINOR \)],
55+
[AC_MSG_WARN([Minimum required CUDA version for device code: $NVCC_CUDA_MIN_REQUIRED_MAJOR.$NVCC_CUDA_MIN_REQUIRED_MINOR])
56+
NVCC=""
57+
])
58+
59+
AS_IF([test "x$NVCC" != "x"], [
60+
AC_ARG_WITH([nvcc-gencode],
61+
[AS_HELP_STRING([--with-nvcc-gencode=(OPTS)], [Build for specific GPU architectures])],
62+
[],
63+
[with_nvcc_gencode=default])
64+
65+
AS_IF([test "x$with_nvcc_gencode" = "xdefault"],
66+
[AS_CASE([$CUDA_MAJOR_VERSION],
67+
[13],
68+
[# offline compilation support for architectures before '<compute/sm/lto>_75' is discontinued
69+
NVCC_ARCH="${ARCH10_CODE} ${ARCH110_CODE} ${ARCH111_CODE} ${ARCH120_CODE} ${ARCH124_CODE} ${ARCH128_CODE} ${ARCH130_CODE} ${ARCH130_PTX}"],
70+
[12],
71+
[AS_CASE([$CUDA_MINOR_VERSION],
72+
[0|1|2|3],
73+
[NVCC_ARCH="${ARCH9_CODE} ${ARCH10_CODE} ${ARCH110_CODE} ${ARCH111_CODE} ${ARCH120_CODE} ${ARCH120_PTX}"],
74+
[4|5|6|7],
75+
[NVCC_ARCH="${ARCH9_CODE} ${ARCH10_CODE} ${ARCH110_CODE} ${ARCH111_CODE} ${ARCH120_CODE} ${ARCH124_CODE} ${ARCH124_PTX}"],
76+
[*],
77+
[NVCC_ARCH="${ARCH9_CODE} ${ARCH10_CODE} ${ARCH110_CODE} ${ARCH111_CODE} ${ARCH120_CODE} ${ARCH124_CODE} ${ARCH128_CODE} ${ARCH128_PTX}"])],
78+
79+
[11],
80+
[AS_CASE([$CUDA_MINOR_VERSION],
81+
[0],
82+
[NVCC_ARCH="${ARCH9_CODE} ${ARCH10_CODE} ${ARCH110_CODE} ${ARCH110_PTX}"],
83+
[*],
84+
[NVCC_ARCH="${ARCH9_CODE} ${ARCH10_CODE} ${ARCH110_CODE} ${ARCH111_CODE} ${ARCH111_PTX}"])],
85+
[*],
86+
[NVCC_ARCH="${ARCH9_CODE} ${ARCH9_PTX}"])],
87+
[NVCC_ARCH="$with_nvcc_gencode"])
88+
BASE_NVCCFLAGS="$BASE_NVCCFLAGS $NVCC_ARCH"
89+
AC_MSG_CHECKING([$NVCC needs explicit c++11 option])
90+
AC_LANG_PUSH([CUDA])
91+
AC_COMPILE_IFELSE([AC_LANG_SOURCE([[
92+
#if __cplusplus < 201103L
93+
#error missing C++11
94+
#endif
95+
]])],
96+
[AC_MSG_RESULT([no])],
97+
[AC_MSG_RESULT([yes])
98+
BASE_NVCCFLAGS="$BASE_NVCCFLAGS -std=c++11"])
99+
AC_LANG_POP
100+
101+
AC_MSG_CHECKING([$NVCC can compile])
102+
AC_LANG_PUSH([CUDA])
103+
AC_COMPILE_IFELSE([AC_LANG_SOURCE([[
104+
#include <cuda_runtime.h>
105+
__global__ void my_kernel(void) {}
106+
int main(void) { my_kernel<<<1, 1>>>(); return 0; }
107+
]])],
108+
[AC_MSG_RESULT([yes])],
109+
[AC_MSG_RESULT([no])
110+
NVCC=""])
111+
AC_LANG_POP
112+
])
113+
])
58114
59115
AM_CONDITIONAL([HAVE_NVCC], [test "x$NVCC" != x])
60116
])

0 commit comments

Comments
 (0)