 
 __global__ void KernelKnapsack(unsigned int *total_buy_options, unsigned int *best_buy_options)
 {
-    __shared__ unsigned int tmp_best_buy_options[NUM_THREADS * 2];
-    unsigned int stride;
+    __shared__ unsigned int tmp_best_buy_options[NUM_BUY_OPTIONS * ELEMENTS_PER_BUY_OPTION];
 
     unsigned int thread_id = threadIdx.x;
-    unsigned int thread_product = blockIdx.x * blockDim.x + threadIdx.x * ELEMENTS_PER_BUY_OPTION;
-    tmp_best_buy_options[thread_id + STORE_ID_OFFSET] = total_buy_options[thread_product + STORE_ID_OFFSET];
-    tmp_best_buy_options[thread_id + PRICE_OFFSET] = total_buy_options[thread_product + PRICE_OFFSET];
+    unsigned int shared_thread_buy_option = thread_id * ELEMENTS_PER_BUY_OPTION;
+    unsigned int global_thread_buy_option = ( blockIdx.x * blockDim.x + thread_id ) * ELEMENTS_PER_BUY_OPTION;
+
+    tmp_best_buy_options[shared_thread_buy_option + STORE_ID_OFFSET] = total_buy_options[global_thread_buy_option + STORE_ID_OFFSET];
+    tmp_best_buy_options[shared_thread_buy_option + PRICE_OFFSET] = total_buy_options[global_thread_buy_option + PRICE_OFFSET];
     __syncthreads();
 
-    for (stride = 1; stride < blockDim.x; stride *= 2)
-    {
-        if (thread_id % (2 * stride) == 0)
+    for (unsigned int stride = 2; stride < blockDim.x; stride *= 2)
+    {
+        if (thread_id % stride == 0)
         {
-            unsigned int next_buy_option_position = thread_id + stride * ELEMENTS_PER_BUY_OPTION;
-
-            if (tmp_best_buy_options[thread_id + PRICE_OFFSET] > tmp_best_buy_options[next_buy_option_position + PRICE_OFFSET])
+            unsigned int next_buy_option_position = shared_thread_buy_option + stride;
+
+            if (tmp_best_buy_options[shared_thread_buy_option + PRICE_OFFSET] > tmp_best_buy_options[next_buy_option_position + PRICE_OFFSET])
             {
-                tmp_best_buy_options[thread_id + STORE_ID_OFFSET] = tmp_best_buy_options[next_buy_option_position + STORE_ID_OFFSET];
-                tmp_best_buy_options[thread_id + PRICE_OFFSET] = tmp_best_buy_options[next_buy_option_position + PRICE_OFFSET];
+                tmp_best_buy_options[shared_thread_buy_option + STORE_ID_OFFSET] = tmp_best_buy_options[next_buy_option_position + STORE_ID_OFFSET];
+                tmp_best_buy_options[shared_thread_buy_option + PRICE_OFFSET] = tmp_best_buy_options[next_buy_option_position + PRICE_OFFSET];
             }
         }
         __syncthreads();

    cudaEventDestroy(start);
    cudaEventDestroy(stop);

    if (IS_DEBUG_ENABLED)
    {
        printAllProductsAllBuyOptions(host_all_products_buy_options);
    }

    printBestBuyOptions(best_buy_options);

    if ( areResultsValid( host_all_products_buy_options, best_buy_options ) )
    {

   unsigned int *tmp_best_buy_options = (unsigned int *) malloc( NUM_PRODUCTS * ELEMENTS_PER_BUY_OPTION * sizeof(unsigned int) );

   getBestBuyOptions(all_products_buy_options, tmp_best_buy_options);

   if (DEBUG_LEVEL >= 1)
   {
       printBestBuyOptions(tmp_best_buy_options);
   }

   for (unsigned int product_iteration = 0; product_iteration < NUM_PRODUCTS * ELEMENTS_PER_BUY_OPTION; product_iteration += ELEMENTS_PER_BUY_OPTION)
   {
