From 65c9d0073066a9f41ba902e037b4a3fe1e83bb01 Mon Sep 17 00:00:00 2001
From: Victor Julien <victor@inliniac.net>
Date: Thu, 15 Apr 2010 23:50:16 +0200
Subject: [PATCH] Remove duplicate cuda kernel file.

---
 src/suricata_kernel.cu | 84 ------------------------------------------
 1 file changed, 84 deletions(-)
 delete mode 100644 src/suricata_kernel.cu

diff --git a/src/suricata_kernel.cu b/src/suricata_kernel.cu
deleted file mode 100644
index 305f2a1978..0000000000
--- a/src/suricata_kernel.cu
+++ /dev/null
@@ -1,84 +0,0 @@
-/**
- * Copyright (c) 2010 Open Information Security Foundation.
- *
- * \author Anoop Saldanha <poonaatsoc@gmail.com>
- *
- * \file The Cuda kernel for MPM B2G.
- *
- * \todo This is a basic version of the kernel.  Modify it to support multiple
- *       blocks of threads.  Make use of shared memory/texture memory.
- */
-
-#define B2G_CUDA_Q 2
-#define CUDA_THREADS 16
-#define B2G_CUDA_HASHSHIFT 4
-#define B2G_CUDA_TYPE unsigned int
-#define B2G_CUDA_HASH16(a, b) (((a) << B2G_CUDA_HASHSHIFT) | (b))
-#define u8_tolower(c) g_u8_lowercasetable[(c)]
-
-extern "C"
-__global__ void B2gCudaSearchBNDMq(unsigned int *offsets,
-                                   unsigned int *B2G,
-                                   unsigned char *g_u8_lowercasetable,
-                                   unsigned char *buf,
-                                   unsigned short arg_buflen,
-                                   unsigned int m)
-{
-    unsigned int pos = m - B2G_CUDA_Q + 1;
-    B2G_CUDA_TYPE d;
-    unsigned short h;
-    unsigned int j;
-    unsigned int first;
-    unsigned int tid = threadIdx.x;
-    unsigned short tid_chunk = arg_buflen / CUDA_THREADS;
-    unsigned short jump;
-    unsigned short buflen;
-
-    if (tid_chunk < m)
-        tid_chunk = m;
-
-    jump = tid_chunk * tid;
-    if ((jump + tid_chunk) > arg_buflen)
-        return;
-
-    buflen = tid_chunk * 2 - 1;
-    if ((tid == CUDA_THREADS - 1) || ((jump + buflen) > arg_buflen)) {
-        buflen = arg_buflen - jump;
-    }
-
-    j = 0;
-    while (j < buflen) {
-        offsets[jump + j] = 0;
-        j++;
-    }
-
-    while (pos <= (buflen - B2G_CUDA_Q + 1)) {
-        h = B2G_CUDA_HASH16(u8_tolower(buf[jump + pos - 1]), u8_tolower(buf[jump + pos]));
-        d = B2G[h];
-
-        if (d != 0) {
-            j = pos;
-            first = pos - (m - B2G_CUDA_Q + 1);
-
-            do {
-                j = j - 1;
-                if (d >= (1 << (m - 1))) {
-                    if (j > first) {
-                        pos = j;
-                    } else {
-                        offsets[j + jump] = 1;
-                    }
-                }
-
-                if (j == 0)
-                    break;
-
-                h = B2G_CUDA_HASH16(u8_tolower(buf[jump + j - 1]), u8_tolower(buf[jump + j]));
-                d = (d << 1) & B2G[h];
-            } while (d != 0);
-        }
-        pos = pos + m - B2G_CUDA_Q + 1;
-    }
-
-    return;
-}