sm_52 compilation error, test for cuda version > 11.6

[torehl]

Hi,

 when compiling ucc 1.2.0 using cuda-12.2.1 toolkit (and ROCm 5.6.0), compilation failed for sm_52.  Maxwell was deprecated from Cuda 11.6 or newer.

 Since I have Pascal or newer hardware at our lab, I uncommented it, and ucc 1.2.0 built. Probably need a test for this.

`torel@srl-login1:~/workspace/UCC/ucc-1.2.0$ diff -auT configure.ORIG configure
--- configure.ORIG 2023-11-01 12:03:35.000000000 +0100
+++ configure 2023-11-01 11:54:05.000000000 +0100
@@ -24218,7 +24218,7 @@
CUDA_MIN_REQUIRED_MAJOR=11
CUDA_MIN_REQUIRED_MINOR=0

• ARCH7="-gencode=arch=compute_52,code=sm_52"

• ARCH7="-gencode=arch=compute_52,code=sm_52"

  ARCH8="-gencode=arch=compute_60,code=sm_60
  -gencode=arch=compute_61,code=sm_61
  -gencode=arch=compute_61,code=compute_61"`

Ref.: https://arnon.dk/matching-sm-architectures-arch-and-gencode-for-various-nvidia-cards/

`/bin/bash ../../../../../../cuda_lt.sh "/bin/sh ../../../../../libtool" ec_cuda_executor.lo /cm/shared/apps/cuda12.2/toolkit/12.2.1/bin/nvcc -c ../../../../../../src/components/ec/cuda/kernel/ec_cuda_executor.cu -O3 -g -DNDEBUG -I/home/torel/workspace/UCC/ucc-1.2.0/Build-x86_64 -I/home/torel/workspace/UCC/ucc-1.2.0/src -I/home/torel/workspace/UCC/ucc-1.2.0/Build-x86_64/src --compiler-options -fno-rtti,-fno-exceptions -gencode=arch=compute_52,code=sm_52 -gencode=arch=compute_60,code=sm_60 -gencode=arch=compute_61,code=sm_61 -gencode=arch=compute_61,code=compute_61 -gencode=arch=compute_70,code=sm_70 -gencode=arch=compute_70,code=compute_70 -gencode=arch=compute_75,code=sm_75 -gencode=arch=compute_80,code=sm_80 -gencode=arch=compute_80,code=compute_80 -gencode=arch=compute_86,code=sm_86 -gencode=arch=compute_86,code=compute_86 -gencode=arch=compute_90,code=sm_90 -gencode=arch=compute_90,code=compute_90
/cm/shared/apps/cuda12.2/toolkit/12.2.1/bin/nvcc -c ../../../../../../src/components/ec/cuda/kernel/ec_cuda_executor.cu -O3 -g -DNDEBUG -I/home/torel/workspace/UCC/ucc-1.2.0/Build-x86_64 -I/home/torel/workspace/UCC/ucc-1.2.0/src -I/home/torel/workspace/UCC/ucc-1.2.0/Build-x86_64/src --compiler-options -fno-rtti,-fno-exceptions -gencode=arch=compute_52,code=sm_52 -gencode=arch=compute_60,code=sm_60 -gencode=arch=compute_61,code=sm_61 -gencode=arch=compute_61,code=compute_61 -gencode=arch=compute_70,code=sm_70 -gencode=arch=compute_70,code=compute_70 -gencode=arch=compute_75,code=sm_75 -gencode=arch=compute_80,code=sm_80 -gencode=arch=compute_80,code=compute_80 -gencode=arch=compute_86,code=sm_86 -gencode=arch=compute_86,code=compute_86 -gencode=arch=compute_90,code=sm_90 -gencode=arch=compute_90,code=compute_90 -Xcompiler -fPIC -o ./.libs/ec_cuda_executor.o
../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(30): error: function "operator+(const __half &, const __half &)" has already been defined
attribute((device)) inline attribute((always_inline)) __half operator+(const __half &lh, const __half &rh)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(34): error: function "operator-(const __half &, const __half &)" has already been defined
attribute((device)) inline attribute((always_inline)) __half operator-(const __half &lh, const __half &rh)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(38): error: function "operator*(const __half &, const __half &)" has already been defined
attribute((device)) inline attribute((always_inline)) __half operator*(const __half &lh, const __half &rh)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(42): error: function "operator/(const __half &, const __half &)" has already been defined
attribute((device)) inline attribute((always_inline)) __half operator/(const __half &lh, const __half &rh)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(47): error: function "operator+=(__half &, const __half &)" has already been defined
attribute((device)) inline attribute((always_inline)) __half &operator+=(__half &lh, const __half &rh)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(52): error: function "operator-=(__half &, const __half &)" has already been defined
attribute((device)) inline attribute((always_inline)) __half &operator-=(__half &lh, const __half &rh)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(57): error: function "operator*=(__half &, const __half &)" has already been defined
attribute((device)) inline attribute((always_inline)) __half &operator*=(__half &lh, const __half &rh)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(62): error: function "operator/=(__half &, const __half &)" has already been defined
attribute((device)) inline attribute((always_inline)) __half &operator/=(__half &lh, const __half &rh)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(69): error: function "operator++(__half &)" has already been defined
attribute((device)) inline attribute((always_inline)) __half &operator++(__half &h)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(76): error: function "operator--(__half &)" has already been defined
attribute((device)) inline attribute((always_inline)) __half &operator--(__half &h)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(83): error: function "operator++(__half &, int)" has already been defined
attribute((device)) inline attribute((always_inline)) __half operator++(__half &h, const int ignored)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(94): error: function "operator--(__half &, int)" has already been defined
attribute((device)) inline attribute((always_inline)) __half operator--(__half &h, const int ignored)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(107): error: function "operator+(const __half &)" has already been defined
attribute((device)) inline attribute((always_inline)) __half operator+(const __half &h)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(111): error: function "operator-(const __half &)" has already been defined
attribute((device)) inline attribute((always_inline)) __half operator-(const __half &h)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(117): error: function "operator==(const __half &, const __half &)" has already been defined
attribute((device)) inline attribute((always_inline)) bool operator==(const __half &lh, const __half &rh)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(121): error: function "operator!=(const __half &, const __half &)" has already been defined
attribute((device)) inline attribute((always_inline)) bool operator!=(const __half &lh, const __half &rh)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(125): error: function "operator>(const __half &, const __half &)" has already been defined
attribute((device)) inline attribute((always_inline)) bool operator>(const __half &lh, const __half &rh)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(129): error: function "operator<(const __half &, const __half &)" has already been defined
attribute((device)) inline attribute((always_inline)) bool operator<(const __half &lh, const __half &rh)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(133): error: function "operator>=(const __half &, const __half &)" has already been defined
attribute((device)) inline attribute((always_inline)) bool operator>=(const __half &lh, const __half &rh)
^

../../../../../../src/components/ec/cuda/kernel/ec_cuda_half_sm52.h(137): error: function "operator<=(const __half &, const __half &)" has already been defined
attribute((device)) inline attribute((always_inline)) bool operator<=(const __half &lh, const __half &rh)
^

20 errors detected in the compilation of "../../../../../../src/components/ec/cuda/kernel/ec_cuda_executor.cu".
make[4]: *** [Makefile:680: ec_cuda_executor.lo] Error 1
make[4]: Leaving directory '/global/D1/homes/torel/workspace/UCC/ucc-1.2.0/Build-x86_64/src/components/ec/cuda/kernel'
make[3]: *** [Makefile:700: all-recursive] Error 1
make[3]: Leaving directory '/global/D1/homes/torel/workspace/UCC/ucc-1.2.0/Build-x86_64/src/components/ec/cuda'
make[2]: *** [Makefile:1399: all-recursive] Error 1
make[2]: Leaving directory '/global/D1/homes/torel/workspace/UCC/ucc-1.2.0/Build-x86_64/src'
make[1]: *** [Makefile:568: all-recursive] Error 1
make[1]: Leaving directory '/global/D1/homes/torel/workspace/UCC/ucc-1.2.0/Build-x86_64'
make: *** [Makefile:494: all] Error 2
`

[Sergei-Lebedev]

Hi @torehl, I don't think this is because of deprecation of Maxwell, but rather due to the issue with half operators. It was fixed by this PR #847