From ee11b78056d33756d9d2f0933f90d41a6f3cf3a7 Mon Sep 17 00:00:00 2001
From: "nicolas.marie" <nicolas.marie@ensiie.eu>
Date: Thu, 21 Dec 2023 20:06:41 +0100
Subject: [PATCH] remove useless condition in kernel
---
Projet/CODE/apm/Makefile | 3 +++
Projet/CODE/apm/src/apm_gpu.cu | 12 ++++--------
2 files changed, 7 insertions(+), 8 deletions(-)
diff --git a/Projet/CODE/apm/Makefile b/Projet/CODE/apm/Makefile
index a605f6d..2b3cde1 100644
--- a/Projet/CODE/apm/Makefile
+++ b/Projet/CODE/apm/Makefile
@@ -27,6 +27,9 @@ test: all
./apm 0 dna/test.fa test123
./apm_omp 0 dna/test.fa test123
./apm_gpu 0 dna/test.fa test123
+ ./apm 0 dna/xenoRefMrna.fa ttggaa
+ ./apm_omp 0 dna/xenoRefMrna.fa ttggaa
+ ./apm_gpu 0 dna/xenoRefMrna.fa ttggaa
.PHONY: clean
clean:
diff --git a/Projet/CODE/apm/src/apm_gpu.cu b/Projet/CODE/apm/src/apm_gpu.cu
index a9eb4fe..4c070a1 100644
--- a/Projet/CODE/apm/src/apm_gpu.cu
+++ b/Projet/CODE/apm/src/apm_gpu.cu
@@ -210,15 +210,10 @@ levenshtein_cu(char *find, char *buf, int len, int n_bytes,
}
}
- int res = 0;
-
if (column[len] <= approx_factor)
{
- res = 1;
+ atomicAdd(result, 1);
}
-
- //{result_vec[tId] = 1;}//its a match
- atomicAdd(result, res);
}
int
@@ -339,7 +334,7 @@ main(int argc, char **argv)
int maxsharedmem;
cudaDeviceGetAttribute(&maxsharedmem, cudaDevAttrMaxSharedMemoryPerBlock, 0);
int NTBB = MIN2(maxsharedmem / ((size_pattern + 1) * sizeof(int)),
- NUMBER_THREADS_BY_BLOCK);
+ NUMBER_THREADS_BY_BLOCK);
int Nshared = (size_pattern + 1) * NTBB * sizeof(int);
#if APM_DEBUG
@@ -350,6 +345,7 @@ main(int argc, char **argv)
#endif
off_t offset = 0;
+
while (offset < filesize)
{
#if APM_DEBUG
@@ -383,7 +379,7 @@ main(int argc, char **argv)
//}
levenshtein_cu <<< NB, NTBB, Nshared>>>(
- pattern_dev, buf_dev, size_pattern,
+ pattern_dev, buf_dev, size_pattern,
n_bytes, approx_factor, result_dev);
__cudaCheckErrors("Kernel execution of levenshtein_cu failed");
--
GitLab