                                           const int filterSize, const int imgSizeY, const int imgSizeX, const int paddingStart, const int moduleStride,
                                           const int numImgColors, const int numGroups,
                                           const float scaleTargets, const float scaleOutputs) {
-    // avoid bank conflict by re-organizing the data structure, and improve band width by using 'float2' instead of 'float'
-    __shared__ float2 shFilters[filterCacheF/2][colorsPerThread*B_Y/2 + 2];
-    __shared__ float2 shHidActs[filterCacheH][B_X*imgsPerThread/2];
-
-    const int tx = threadIdx.x, ty = threadIdx.y;
+    __shared__ float shFilters[colorsPerThread*B_Y][filterCacheF];
+    __shared__ float shHidActs[filterCacheH][B_X*imgsPerThread];
+
     const int numImgBlocks = DIVUP(numImages,B_X*imgsPerThread);
-    //const int blockCaseIdx = (blockIdx.x % numImgBlocks) * B_X*imgsPerThread;
-    //const int myCaseIdx = blockCaseIdx + threadIdx.x;
-
-    //const int imgColorIdx = (blockIdx.x / numImgBlocks) * B_Y*colorsPerThread; // color idx globally
+    const int blockCaseIdx = (blockIdx.x % numImgBlocks) * B_X*imgsPerThread;
+    const int myCaseIdx = blockCaseIdx + threadIdx.x;
+
+    const int imgColorIdx = (blockIdx.x / numImgBlocks) * B_Y*colorsPerThread; // color idx globally
     const int numFilterColors = numImgColors / numGroups;
-    const int blockGroupIdx = (blockIdx.x / numImgBlocks) * B_Y * colorsPerThread / numFilterColors;
-    const int filterColorIdx = ((blockIdx.x / numImgBlocks) * B_Y * colorsPerThread) % numFilterColors; // color idx within group
+    const int blockGroupIdx = imgColorIdx / numFilterColors;
+    const int filterColorIdx = imgColorIdx % numFilterColors; // color idx within group
     const int numFiltersPerGroup = numFilters / numGroups;
-    //const int blockFilterIdx = blockGroupIdx * numFiltersPerGroup;
-
-    //const int blockPixelIdx = blockIdx.y;
-    const int blockPixelIdxX = blockIdx.y % imgSizeX;
-    const int blockPixelIdxY = blockIdx.y / imgSizeX;
-
-    //const int filterPixels = filterSize * filterSize;
-    const int filterPxlAll = filterSize * filterSize * numFilters;
-    //const int imgPixels = imgSizeY * imgSizeX;
-    //const int tidx = threadIdx.y * B_X + threadIdx.x;
+    const int blockFilterIdx = blockGroupIdx * numFiltersPerGroup;
+
+    const int blockPixelIdx = blockIdx.y;
+    const int blockPixelIdxX = blockPixelIdx % imgSizeX;
+    const int blockPixelIdxY = blockPixelIdx / imgSizeX;
+
+    const int filterPixels = filterSize * filterSize;
+    const int imgPixels = imgSizeY * imgSizeX;
+    const int tidx = threadIdx.y * B_X + threadIdx.x;
 //    const int hidActLoadY = threadIdx.y % B_Y, hidActLoadX = threadIdx.x % B_X;
     //const int hidActLoadY = tidx / (B_X*imgsPerThread), hidActLoadX = tidx % (B_X*imgsPerThread);
-    //const int filtersLoadY = tidx / filterCacheF, filtersLoadX = tidx % filterCacheF;
+    const int filtersLoadY = tidx / filterCacheF, filtersLoadX = tidx % filterCacheF;
     // nvcc is behaving idiotically again, these useless declarations save registers
     //const int outputY = threadIdx.y, outputX = threadIdx.x;
     //const int ty = threadIdx.y, tx = threadIdx.x;
-    //const int numModules = numModulesY * numModulesX;
-    const int numImgMods = numModulesY * numModulesX * numImages;  
-    const int hidActsOffset = (blockGroupIdx * numFiltersPerGroup + ty) * numImgMods + (blockIdx.x % numImgBlocks) * B_X * imgsPerThread + tx;
-    const int filtersOffset = blockGroupIdx * numFiltersPerGroup + (filterColorIdx + (ty * B_X + tx) / filterCacheH) * filterPxlAll + (ty * B_X + tx) % filterCacheH;
+    const int numModules = numModulesY * numModulesX;
+    const int hidActsOffset = (blockFilterIdx + threadIdx.y) * numImages * numModules + myCaseIdx;
+    const int filtersOffset = blockFilterIdx + (filterColorIdx + filtersLoadY) * filterPixels * numFilters + filtersLoadX;
 //    hidActs += (blockFilterIdx + threadIdx.y) * numImages * numModules + myCaseIdx;
 //    filters += blockFilterIdx + (filterColorIdx + filtersLoadY) * filterPixels * numFilters + filtersLoadX;
-    targets += ((blockIdx.x / numImgBlocks) * B_Y * colorsPerThread + ty) * imgSizeY * imgSizeX * numImages + blockIdx.y * numImages + (blockIdx.x % numImgBlocks) * 128 + tx;
+    targets += (imgColorIdx + threadIdx.y) * imgPixels * numImages + blockPixelIdx * numImages + myCaseIdx;
 
     float prod[colorsPerThread][imgsPerThread];
     #pragma unroll

                                conv_img_acts_manycolor_preloadfh_ty_8_tx_32_c_8_ff_32_fh_16_tex< 8, 32, 4, 8, 32, 16, false, false, true ><<<blocks, threads, 0, stream>>>(hidActs.getTextureObject(), filters.getTextureObject(), targets.getDevData(), numModulesY, numModulesX, numImages, numFilters, filterSize, imgSizeY, imgSizeX, paddingStart, moduleStride, numImgColors, numGroups, scaleTargets, scaleOutput);
                            }
                            else if (numImages % 64 == 0) {
                                cudaFuncSetCacheConfig(conv_img_acts_manycolor_kepler_f_vec_h_vec_ff_32_fh_16 < 8, 32, 2, 8, 32, 16, false, false, true >, cudaFuncCachePreferShared);
                                conv_img_acts_manycolor_kepler_f_vec_h_vec_ff_32_fh_16 < 8, 32, 2, 8, 32, 16, false, false, true ><<<blocks, threads, 0, stream>>>(hidActs.getDevData(), filters.getDevData(), targets.getDevData(), numModulesY, numModulesX, numImages, numFilters, filterSize, imgSizeY, imgSizeX, paddingStart, moduleStride, numImgColors, numGroups, scaleTargets, scaleOutput);
                            }
                            else if (numImages % 32 == 0) {
                                cudaFuncSetCacheConfig(conv_img_acts_manycolor_kepler_f_vec_h_vec_ff_32_fh_16 < 8, 32, 1, 8, 32, 16, false, false, true >, cudaFuncCachePreferShared);
                                conv_img_acts_manycolor_kepler_f_vec_h_vec_ff_32_fh_16 < 8, 32, 1, 8, 32, 16, false, false, true ><<<blocks, threads, 0, stream>>>(hidActs.getDevData(), filters.getDevData(), targets.getDevData(), numModulesY, numModulesX, numImages, numFilters, filterSize, imgSizeY, imgSizeX, paddingStart, moduleStride, numImgColors, numGroups, scaleTargets, scaleOutput);
                            }
                            else if (numImages % 16 == 0) {
                                cudaFuncSetCacheConfig(conv_img_acts_manycolor_kepler < 8, 32, 1, 8, 32, 16, false, false, true >, cudaFuncCachePreferShared);

                                conv_img_acts_manycolor_preloadfh_ty_8_tx_32_c_8_ff_32_fh_16_tex< 8, 32, 4, 8, 32, 16, true, false, true ><<<blocks, threads, 0, stream>>>(hidActs.getTextureObject(), filters.getTextureObject(), targets.getDevData(), numModulesY, numModulesX, numImages, numFilters, filterSize, imgSizeY, imgSizeX, paddingStart, moduleStride, numImgColors, numGroups, scaleTargets, scaleOutput);
                            }
                            else if (numImages % 64 == 0) {
                                cudaFuncSetCacheConfig(conv_img_acts_manycolor_kepler < 8, 32, 2, 8, 32, 16, true, false, true >, cudaFuncCachePreferShared);
                                conv_img_acts_manycolor_kepler < 8, 32, 2, 8, 32, 16, true, false, true ><<<blocks, threads, 0, stream>>>(hidActs.getDevData(), filters.getDevData(), targets.getDevData(), numModulesY, numModulesX, numImages, numFilters, filterSize, imgSizeY, imgSizeX, paddingStart, moduleStride, numImgColors, numGroups, scaleTargets, scaleOutput);
                            }
                            else if (numImages % 32 == 0) {
                                cudaFuncSetCacheConfig(conv_img_acts_manycolor_kepler < 8, 32, 1, 8, 32, 16, true, false, true >, cudaFuncCachePreferShared);
                                conv_img_acts_manycolor_kepler < 8, 32, 1, 8, 32, 16, true, false, true ><<<blocks, threads, 0, stream>>>(hidActs.getDevData(), filters.getDevData(), targets.getDevData(), numModulesY, numModulesX, numImages, numFilters, filterSize, imgSizeY, imgSizeX, paddingStart, moduleStride, numImgColors, numGroups, scaleTargets, scaleOutput);
                            }
                            else if (numImages % 16 == 0) {
                                cudaFuncSetCacheConfig(conv_img_acts_manycolor_kepler < 8, 32, 1, 8, 32, 16, true, false, true >, cudaFuncCachePreferShared);
