fixed compilation issues with latest cuda

Author: 20177282

.github/workflows/test-build.yml

@@ -24,4 +24,4 @@ jobs:
           sudo apt-get update
           sudo apt-get -y install cuda-toolkit-12-8
     - name: make check
-      run: bash install.sh -c
+      run: bash install.sh -c -g

VERSION

@@ -1,2 +1,2 @@
 cpu:3.2.5
-gpu:3.4.3
+gpu:3.4.4

install.sh

@@ -32,7 +32,7 @@ $ch	-s or --statistics    enable costly statistics (may impact runtime)
 $ch	-a or --all           enable all above flags except 'assert' and 'verbosity'
 $ch	--ncolors             disable colors in all solver outputs
 $ch	--clean=<target>      remove old installation of <cpu | gpu | all> solvers
-$ch	--standard=<n>        compile with <11 | 14 | 17> c++ standard
+$ch	--standard=<n>        compile with <17 | 20> c++ standard
 $ch	--gextra="flags"      pass extra "flags" to the GPU compiler (nvcc)
 $ch --cextra="flags"      pass extra "flags" to the CPU compiler (g++)
 EOF
@@ -116,7 +116,7 @@ if [[ "$clean" != "" ]] && [[ "$clean" != "cpu" ]] && [[ "$clean" != "gpu" ]] &&
 	error "invalid clean target '$clean'"
 fi
 
-if [ ! $standard = 14 ] && [ ! $standard = 17 ]; then 
+if [ ! $standard = 17 ] && [ ! $standard = 20 ]; then 
 	error "invalid c++ standard '$standard'"
 fi
 
@@ -397,6 +397,7 @@ if [[ $pedantic = 1 ]]; then log "  turning off 'pedantic' due to incompatibilit
 
 # default flags
 INCLUDE="../../dep"
+RELOC="-rdc=true"
 EXTLAMBDA="--expt-extended-lambda"
 RELAXEDEXPR="--expt-relaxed-constexpr"
 OPTIMIZE="-O3"
@@ -433,13 +434,21 @@ if [ -d dep ]; then
 	NVCCFLAGS="$NVCCFLAGS -I$INCLUDE"
 fi
 
+NVCCFLAGS="$NVCCFLAGS $RELOC"
+
 if [[ $gextra != "" ]]; then NVCCFLAGS="$NVCCFLAGS $gextra"; fi
 if [[ $cextra != "" ]]; then CCFLAGS="$CCFLAGS $cextra"; fi
 
+first=$(printf "%s" "$NVCCFLAGS" | cut -d' ' -f1-5)
+second=$(printf "%s" "$NVCCFLAGS" | cut -d' ' -f6-8)
+rest=$(printf "%s" "$NVCCFLAGS" | cut -d' ' -f9-)
+
 log ""
 log "building with:"
-log "'$NVCCFLAGS"
-log "-Xcompiler $CCFLAGS'"
+[ ! -z "$first" ] && log "$(printf " %s\n" "$first")"
+[ ! -z "$second" ] && log "$(printf " %s\n" "$second")"
+[ ! -z "$rest" ] && log "$(printf " %s\n" "$rest")"
+log " -Xcompiler $CCFLAGS'"
 log ""
 
 [ ! -f $gputemplate ] && error "cannot find the GPU makefile template"

src/cpu/version.h

@@ -26,3 +26,6 @@ namespace ParaFROST {
 	const char* date();
 
 }
+#define VERSION "3.2.5"
+#define OSYSTEM "linux rig-muhos 5.15.167.4-microsoft-standard-wsl2 x86_64"
+#define DATE "Sun Apr 27 21:41:37 CEST 2025"

src/gpu/atomics.cuh

@@ -25,26 +25,8 @@ along with this program.  If not, see <https://www.gnu.org/licenses/>.
 
 namespace ParaFROST {
 
-	template<class T>
-	_PFROST_D_ T atomicAggInc(T* counter) {
-		const uint32 mask = __activemask(), total = __popc(mask);
-		uint32 laneMask;
-		laneMask_lt(laneMask);
-		const T prefix = (T)__popc(mask & laneMask);
-		const int lowest_lane = __ffs(mask) - 1;
-		T warpRes = prefix ? 0 : atomicAdd(counter, total);
-		warpRes = __shfl_sync(mask, warpRes, lowest_lane);
-		return (prefix + warpRes);
-	}
-
-	template<class T, class R>
-	_PFROST_D_ void atomicAggMax(T* counter, const R ref) {
-		const uint32 mask = __activemask(), max_id = (32 - __clz(mask)) - 1;
-		uint32 lane_id;
-		laneId(lane_id);
-		if (lane_id == max_id)
-			atomicMax(counter, ref);
-	}
+
+	
 
 }
 

src/gpu/bounded.cuh

@@ -25,6 +25,7 @@ along with this program.  If not, see <https://www.gnu.org/licenses/>.
 #include "function.cuh"
 #include "ifthenelse.cuh"
 #include "equivalence.cuh"
+#include "shared.cuh"
 
 namespace ParaFROST {
 
@@ -65,8 +66,6 @@ namespace ParaFROST {
 
 	//==========================================//
 
-	__device__ int   lastEliminatedID;
-
 	#define ADD_RESOLVENT \
 	{ \
 		int rsize; \
@@ -262,21 +261,7 @@ namespace ParaFROST {
 	//=========================================================//
 	// kernels
 	//=========================================================//
-
-	__global__ void reset_id() { lastEliminatedID = -1; }
-
-	__global__ void print_id() { printf("c lastEliminatedID = %d\n", lastEliminatedID); }
 
-	__global__ void mapfrozen_k(const uint32* __restrict__ frozen, uint32* __restrict__ varcore, const uint32 size)
-	{
-		grid_t tid = global_tx;
-		while (tid < size) {
-			assert(frozen[tid] && frozen[tid] < NOVAR);
-			varcore[frozen[tid]] = tid;
-			tid += stride_x;
-		}
-	}
-
 	// Macros for checking applicability of variable elimination
 	#define MEMORY_SAFE_DBG \
 		if ((addedPos + nAddedCls) > cnf->refs().capacity()) { \
@@ -491,6 +476,19 @@ namespace ParaFROST {
 		}
 	}
 
+	_PFROST_IN_D_ uint32 laneId() {
+		uint32 id;
+		asm("mov.u32 %0, %%laneid;" : "=r"(id));
+		return id;
+	}
+
+	template<class T, class R>
+	_PFROST_IN_D_ void atomicAggMax(T* counter, const R ref) {
+		const uint32 mask = __activemask(), max_id = (32 - __clz(mask)) - 1;
+		if (laneId() == max_id)
+			atomicMax(counter, ref);
+	}
+
 	__global__ void ve_k_2(
 		CNF* __restrict__ cnf,
 		OT* __restrict__ ot,
@@ -551,32 +549,6 @@ namespace ParaFROST {
 		}
 	}
 
-	__global__ void resizeCNF_k(CNF* cnf,
-		const uint32* __restrict__ type,
-		const uint32* __restrict__ rpos,
-		const S_REF* __restrict__ rref,
-		const int verbose)
-	{
-		if (lastEliminatedID >= 0) {
-			const uint32 lastAdded = type[lastEliminatedID];
-			const uint32 lastAddedPos = rpos[lastEliminatedID];
-			const S_REF  lastAddedRef = rref[lastEliminatedID];
-			assert(lastAdded < NOVAR);
-			assert(lastAddedPos < NOVAR);
-			assert(lastAddedRef < GNOREF);
-			assert(RECOVERTYPE(lastAdded) < TYPE_MASK);
-			const uint32 lastAddedCls = RECOVERADDEDCLS(lastAdded);
-			const uint32 lastAddedLits = RECOVERADDEDLITS(lastAdded);
-			assert(lastAddedCls && lastAddedCls <= ADDEDCLS_MAX);
-			assert(lastAddedLits && lastAddedLits <= ADDEDLITS_MAX);
-			const S_REF lastAddedBuckets = lastAddedLits + DC_NBUCKETS * lastAddedCls;
-			const S_REF data_size = lastAddedBuckets + lastAddedRef;
-			const uint32 cs_size = lastAddedCls + lastAddedPos;
-			cnf->resize(data_size, cs_size);
-			if (verbose > 1) printf("c   resized CNF to %d clauses and %lld data for a last ID %d\n", cs_size, data_size, lastEliminatedID);
-		}
-	}
-
 } 
 
 

src/gpu/cache.cuh

@@ -67,8 +67,6 @@ namespace ParaFROST {
 
 	};
 
-	extern CACHER cacher;
-
 }
 
 #endif