refs #6, #27: IPv4 router passes "match" test.

* It means that the CPU version and GPU version yields the same results upon the same inputs.
ANLAB-KAIST · Feb 24, 2016 · 2a4050c · 2a4050c
1 parent 171037a
commit 2a4050c
Show file tree

Hide file tree

Showing 2 changed files with 19 additions and 15 deletions.
diff --git a/elements/ip/IPlookup_kernel.cu b/elements/ip/IPlookup_kernel.cu
@@ -40,8 +40,8 @@ __global__ void ipv4_route_lookup_cuda(
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
 
  if (idx < count) {
- uint16_t batch_idx = batch_ids[idx];
- uint16_t item_idx  = item_ids[idx];
+ uint8_t batch_idx = batch_ids[idx];
+ uint16_t item_idx = item_ids[idx];
  struct datablock_kernel_arg *db_dest_addrs = datablocks[dbid_ipv4_dest_addrs_d];
  struct datablock_kernel_arg *db_results = datablocks[dbid_ipv4_lookup_results_d];
  uint32_t daddr = ((uint32_t*) db_dest_addrs->batches[batch_idx].buffer_bases_in)[item_idx];

diff --git a/tests/test_ipv4route.cc b/tests/test_ipv4route.cc
@@ -1,4 +1,6 @@
 #include <cstdint>
+#include <cstdlib>
+#include <cstdio>
 #ifdef USE_CUDA
 #include <cuda_runtime.h>
 #endif
@@ -54,6 +56,7 @@ class IPLookupCUDAMatchTest : public ::testing::TestWithParam<int> {
 protected:
  virtual void SetUp() {
  cudaSetDevice(GetParam());
+ srand(0); // for deterministic nexthop results
  ipv4route::load_rib_from_file(tables, "configs/routing_info.txt");
  tbl24_h = (uint16_t *) malloc(sizeof(uint16_t)
  * ipv4route::get_TBL24_size());
@@ -87,7 +90,7 @@ class IPLookupCUDAMatchTest : public ::testing::TestWithParam<int> {
  void *tbllong_d;
 };
 
-TEST_P(IPLookupCUDAMatchTest, SinglePacket) {
+TEST_P(IPLookupCUDAMatchTest, SingleBatch) {
  void *k = ipv4_route_lookup_get_cuda_kernel();
  const char *ip1 = "118.223.0.3";
  const char *ip2 = "58.29.89.55";
@@ -100,8 +103,9 @@ TEST_P(IPLookupCUDAMatchTest, SinglePacket) {
  EXPECT_NE(0, cpu_results[0]);
  EXPECT_NE(0, cpu_results[1]);
 
- const size_t num_batches = 1;
- const size_t num_pkts = 2;
+ const uint32_t num_batches = 1;
+ const uint32_t num_pkts = 2;
+ const uint32_t count = num_batches * num_pkts;
 
  struct datablock_kernel_arg *datablocks[2];
  const size_t db_arg_size = sizeof(struct datablock_kernel_arg)
@@ -122,8 +126,8 @@ TEST_P(IPLookupCUDAMatchTest, SinglePacket) {
  uint16_t *output_buffer = (uint16_t *) malloc(output_size);
  ASSERT_NE(nullptr, input_buffer);
  ASSERT_NE(nullptr, output_buffer);
- input_buffer[0] = (uint32_t) ntohl(inet_addr(ip1));
- input_buffer[1] = (uint32_t) ntohl(inet_addr(ip2));
+ input_buffer[0] = (uint32_t) inet_addr(ip1); // ntohl is done inside kernels
+ input_buffer[1] = (uint32_t) inet_addr(ip2);
  output_buffer[0] = 0;
  output_buffer[1] = 0;
  void *input_buffer_d = nullptr;
@@ -167,24 +171,24 @@ TEST_P(IPLookupCUDAMatchTest, SinglePacket) {
  db_arg_size, cudaMemcpyHostToDevice));
  void *dbarray_h[2] = { db_ipv4_dest_addrs_d, db_ipv4_lookup_results_d };
  void *dbarray_d = nullptr;
- uint8_t batch_ids[num_batches] = { 0 };
- uint16_t item_ids[num_pkts]  = { 0, 1 };
+ uint8_t batch_ids[count] = { 0, 0 };
+ uint16_t item_ids[count] = { 0, 1 };
  void *batch_ids_d = nullptr;
  void *item_ids_d = nullptr;
  ASSERT_EQ(cudaSuccess, cudaMalloc(&dbarray_d, sizeof(void*) * 2));
- ASSERT_EQ(cudaSuccess, cudaMalloc(&batch_ids_d, sizeof(uint8_t) * num_batches));
- ASSERT_EQ(cudaSuccess, cudaMalloc(&item_ids_d, sizeof(uint16_t) * num_pkts));
+ ASSERT_EQ(cudaSuccess, cudaMalloc(&batch_ids_d, sizeof(uint8_t) * count));
+ ASSERT_EQ(cudaSuccess, cudaMalloc(&item_ids_d, sizeof(uint16_t) * count));
  ASSERT_NE(nullptr, dbarray_d);
  ASSERT_NE(nullptr, batch_ids_d);
  ASSERT_NE(nullptr, item_ids_d);
  ASSERT_EQ(cudaSuccess, cudaMemcpy(dbarray_d, dbarray_h,
  sizeof(void*) * 2,
  cudaMemcpyHostToDevice));
  ASSERT_EQ(cudaSuccess, cudaMemcpy(batch_ids_d, batch_ids,
- sizeof(uint8_t) * num_batches,
+ sizeof(uint8_t) * count,
  cudaMemcpyHostToDevice));
- ASSERT_EQ(cudaSuccess, cudaMemcpy(item_ids_d, batch_ids,
- sizeof(uint16_t) * num_pkts,
+ ASSERT_EQ(cudaSuccess, cudaMemcpy(item_ids_d, item_ids,
+ sizeof(uint16_t) * count,
  cudaMemcpyHostToDevice));
  void *checkbits_d = nullptr;
 
@@ -220,7 +224,7 @@ TEST_P(IPLookupCUDAMatchTest, SinglePacket) {
  ASSERT_EQ(cudaSuccess, cudaFree(dbarray_d));
 }
 
-TEST_P(IPLookupCUDAMatchTest, Batch) {
+TEST_P(IPLookupCUDAMatchTest, MultipleBatches) {
  void *k = ipv4_route_lookup_get_cuda_kernel();
  ASSERT_EQ(cudaSuccess, cudaDeviceSynchronize());
 }