Squashed 'lib/libmarv/' changes from 444c7dfe..ca62422d

ca62422d Don't build DPX kernels for pre-hopper

git-subtree-dir: lib/libmarv
git-subtree-split: ca62422d2332b801fd4c121094bfe859d004ef6d

Author: milot-mirdita

src/pssmkernels_gapless.cuh

@@ -200,6 +200,11 @@ namespace hardcodedzero{
         __grid_constant__ const SequenceLengthT queryLength,
         __grid_constant__ const PSSM_2D_View<ScoreType> strided_PSSM
     ) {
+        if constexpr (std::is_same_v<ScoreType, short2>) {
+            #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 900
+            return;
+            #endif
+        }
         static_assert(numRegs % 4 == 0);
         static_assert(blocksize % group_size == 0);
         __builtin_assume(blockDim.x == blocksize);
@@ -604,6 +609,11 @@ namespace hardcodedzero{
         __grid_constant__ float2* const multiTileTempStorage,
         __grid_constant__ const size_t tempStorageElementsPerGroup
     ) {
+        if constexpr (std::is_same_v<ScoreType, short2>) {
+            #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 900
+            return;
+            #endif
+        }
         static_assert(numRegs % 4 == 0);
         static_assert(blocksize % group_size == 0);
         __builtin_assume(blockDim.x == blocksize);
@@ -1293,6 +1303,11 @@ namespace kernelparamzero{
         __grid_constant__ const PSSM_2D_View<ScoreType> strided_PSSM,
         __grid_constant__ const ScoreType zero
     ) {
+        if constexpr (std::is_same_v<ScoreType, short2>) {
+            #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 900
+            return;
+            #endif
+        }
         static_assert(numRegs % 4 == 0);
         static_assert(blocksize % group_size == 0);
         __builtin_assume(blockDim.x == blocksize);
@@ -1677,6 +1692,11 @@ namespace kernelparamzero{
         __grid_constant__ const size_t tempStorageElementsPerGroup,
         __grid_constant__ const ScoreType zero
     ) {
+        if constexpr (std::is_same_v<ScoreType, short2>) {
+            #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 900
+            return;
+            #endif
+        }
         static_assert(numRegs % 4 == 0);
         static_assert(blocksize % group_size == 0);
         __builtin_assume(blockDim.x == blocksize);

src/pssmkernels_smithwaterman.cuh

@@ -74,6 +74,11 @@ void amino_gpu_localAlignmentKernel_affinegap_floatOrInt_pssm_singletile(
     __grid_constant__ const ScoreType gapopenscore, 
     __grid_constant__ const ScoreType gapextendscore
 ){
+    if constexpr (std::is_same_v<ScoreType, int>) {
+        #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 900
+        return;
+        #endif
+    }
     static_assert(std::is_same_v<ScoreType, float> || std::is_same_v<ScoreType, int>);
 
     static_assert(groupsize >= 4);
@@ -811,6 +816,11 @@ void amino_gpu_localAlignmentKernel_affinegap_floatOrInt_pssm_multitile(
     __grid_constant__ char* const tempStorage,
     __grid_constant__ const size_t tempBytesPerGroup
 ){
+    if constexpr (std::is_same_v<ScoreType, int>) {
+        #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 900
+        return;
+        #endif
+    }
     static_assert(std::is_same_v<ScoreType, float> || std::is_same_v<ScoreType, int>);
 
     static_assert(groupsize >= 4);