  */
 template <int B_Y, int B_X, int filtersPerThread, int colorsPerThread, int preloadCases, bool scale>
 __launch_bounds__(256, 2)
-__global__ void conv_weight_acts_mc_mf_kepler_preload_ty_8_tx_32_f_4_c_8_r_16(cudaTextureObject_t images, cudaTextureObject_t hidActs, float* targets,
+__global__ void conv_weight_acts_mc_mf_kepler_preload_ty_8_tx_32_f_4_c_6_r_32(float* images, float* hidActs, float* targets,
+                                       const int numImages, const int numFilters,
+                                       const int numModulesY, const int numModulesX,
+                                       const int imgSizeY, const int imgSizeX, const int filterSize,
+                                       const int paddingStart, const int moduleStride, const int imgStride,
+                                       const int numImgColors, const int numGroups, const int sumWidth,
+                                       const float scaleTargets, const float scaleOutputs) {
+    __shared__ float shImages[colorsPerThread * B_Y][preloadCases]; // preload preloadCases cases
+    __shared__ float shHidActs[filtersPerThread * B_X][preloadCases + 1]; // preload preloadCases cases of B_X hidacts
+
+    const int tidx = B_X * threadIdx.y + threadIdx.x;
+    const int loadY = tidx / preloadCases, loadX = tidx % preloadCases;
+
+    const int filterPixels = filterSize * filterSize;
+    const int imgPixels = imgSizeY * imgSizeX;
+
+    const int numFilterBlocks = numFilters / (B_X * filtersPerThread);
+    const int blockModuleChunkIdx = blockIdx.x / numFilterBlocks;
+
+    const int numModuleChunksX = DIVUP(numModulesX, sumWidth);
+//    const int numModuleChunksY = DIVUP(numModulesY, sumWidth);
+
+    const int blockModuleChunkX = blockModuleChunkIdx % numModuleChunksX;
+    const int blockModuleChunkY = blockModuleChunkIdx / numModuleChunksX;
+
+    const int blockModuleStartX = blockModuleChunkX * sumWidth;
+    const int blockModuleStartY = blockModuleChunkY * sumWidth;
+
+//    const int moduleIdx = partialSum * outputModuleIdx;
+    const int blockFilterIdx = filtersPerThread * B_X * (blockIdx.x % numFilterBlocks);
+    const int numModules = numModulesY * numModulesX;
+
+    const int numFiltersPerGroup = numFilters / numGroups;
+    const int blockGroupIdx = blockFilterIdx / numFiltersPerGroup;
+    const int numFilterColors = numImgColors / numGroups;
+
+    const int blockPixelOffset = blockIdx.z; // pixel idx in filter
+    const int blockPixelY = blockPixelOffset / filterSize, blockPixelX = blockPixelOffset % filterSize;
+    const int blockFilterColorIdx = blockIdx.y  * B_Y * colorsPerThread;
+    const int imgColorIdx = blockFilterColorIdx + blockGroupIdx * numFilterColors;
+
+    const int imgOffset = (imgColorIdx + loadY) * imgPixels * imgStride + loadX;
+    const int hidActsOffset = blockFilterIdx * numImages * numModules
+            + loadY * numImages * numModules
+            + loadX;
+//    images += (imgColorIdx + loadY) * imgPixels * imgStride + loadX;
+//
+//    hidActs +=
+//             blockFilterIdx * numImages * numModules
+//            + loadY * numImages * numModules
+//            + loadX;
+
+    targets += blockModuleChunkIdx * numFilters * filterPixels * numFilterColors
+            + (blockFilterColorIdx + threadIdx.y) * filterPixels * numFilters
+            + blockPixelOffset * numFilters
+            + blockFilterIdx
+            + threadIdx.x;
+
+    const int mStartX = max(blockModuleStartX, DIVUP(-blockPixelX - paddingStart, moduleStride));
+    const int mStartY = max(blockModuleStartY, DIVUP(-blockPixelY - paddingStart, moduleStride));
+    const int mEndX = min(numModulesX, min(blockModuleStartX + sumWidth, DIVUP(imgSizeX - blockPixelX - paddingStart, moduleStride)));
+    const int mEndY = min(numModulesY, min(blockModuleStartY + sumWidth, DIVUP(imgSizeY - blockPixelY - paddingStart, moduleStride)));
+
+    const bool doWork = mStartY < mEndY && mStartX < mEndX;
+
+    float* shHidActLoad = &shHidActs[loadY][loadX];
+    float* shImgLoad = &shImages[loadY][loadX];
+
+    float imPreload[preloadCases*colorsPerThread/B_X]; // [6]
+    float haPreload[preloadCases*filtersPerThread/B_Y]; // [16]
+
+    float prod[filtersPerThread][colorsPerThread];
+
+    #pragma unroll
+    for (int f = 0; f < filtersPerThread; f++) {
+        #pragma unroll
+        for (int c = 0; c < colorsPerThread; c++) {
+            prod[f][c] = 0;
+        }
+    }
+    int pixIdx, pixIdxNext, m, mNext;
+
+    conv_weight_acts_mc_mf_kepler_preload_ty_8_tx_32_f_4_c_8_r_16_setCoords(
+            mStartY, mStartX, paddingStart, numModulesX, moduleStride,
+            blockPixelY, blockPixelX, imgSizeX, imgStride,
+            pixIdx, m);
+
+    if (doWork) {
+        #pragma unroll
+        for (int y = 0; y < B_Y * colorsPerThread; y += (B_X * B_Y) / preloadCases) {
+            imPreload[y * preloadCases/(B_X * B_Y)] = images[imgOffset + y * imgPixels * imgStride + pixIdx];
+        }
+
+        #pragma unroll
+        for (int y = 0; y < B_X * filtersPerThread; y += (B_X * B_Y) / preloadCases) {
+            haPreload[y * preloadCases / (B_X * B_Y)] = hidActs[hidActsOffset + y * numImages * numModules + m * numImages];
+        }
+    }
+
+    for (int my = mStartY; my < mEndY; my++) {
+        for (int mx = mStartX; mx < mEndX; mx++) {
+            int myNext = my, mxNext = mx;
+            const bool lastModule = my == mEndY - 1 && mx == mEndX - 1;
+
+            if (!lastModule) {
+                mxNext = mx + 1 == mEndX ? mStartX : mx + 1;
+                myNext = my + (mx + 1 == mEndX);
+            }
+
+            conv_weight_acts_mc_mf_kepler_preload_ty_8_tx_32_f_4_c_8_r_16_setCoords(
+                    myNext, mxNext, paddingStart, numModulesX, moduleStride,
+                    blockPixelY, blockPixelX, imgSizeX, imgStride,
+                    pixIdxNext, mNext);
+
+            for (int caseIdx = 0; caseIdx < numImages; caseIdx += preloadCases) {
+                #pragma unroll
+                for (int y = 0; y < B_Y * colorsPerThread; y += (B_X * B_Y) / preloadCases) {
+                    shImgLoad[(y) * preloadCases] = imPreload[y * preloadCases / (B_X * B_Y)];
+                }
+
+                #pragma unroll
+                for (int y = 0; y < B_X * filtersPerThread; y += (B_X * B_Y) / preloadCases) {
+                    shHidActLoad[y * (preloadCases + 1)] = haPreload[y * preloadCases / (B_X * B_Y)];
+                }
+
+                __syncthreads();
+
+//                const float* im = &images[caseIdx + preloadCases + pixIdx];
+//                const float* ha = &hidActs[caseIdx + preloadCases + m * numImages];
+                int imgOffset2 = imgOffset + caseIdx + preloadCases + pixIdx;
+                int hidActsOffset2 = hidActsOffset + caseIdx + preloadCases + m * numImages;
+                if (caseIdx + preloadCases == numImages) {
+                    pixIdx = pixIdxNext;
+                    m = mNext;
+                    imgOffset2 = imgOffset + pixIdxNext;
+                    hidActsOffset2 = hidActsOffset + mNext * numImages;
+                }
+
+                WA_LOOP(0);
+                WA_LOOP(1);
+                WA_LOOP(2);
+                WA_LOOP(3);
+                WA_LOOP(4);
+
+                WA_LOOP(5);
+                WA_IMLOAD_2(0);
+                WA_LOOP(6);
+                WA_IMLOAD_2(1);
+                WA_LOOP(7);
+                WA_IMLOAD_2(2);
+                WA_LOOP(8);
+                WA_IMLOAD_2(3);
+                WA_LOOP(9);
+                WA_IMLOAD_2(4);
+                WA_LOOP(10);
+                WA_IMLOAD_2(5);
+
+                WA_LOOP(11);
+                WA_HALOAD_2(0);
+                WA_LOOP(12);
+                WA_HALOAD_2(1);
+                WA_LOOP(13);
+                WA_HALOAD_2(2);
+                WA_LOOP(14);
+                WA_HALOAD_2(3);
+                WA_LOOP(15);
+                WA_HALOAD_2(4);
+                WA_LOOP(16);
+                WA_HALOAD_2(5);
+                WA_LOOP(17);
+                WA_HALOAD_2(6);
+                WA_LOOP(18);
+                WA_HALOAD_2(7);
+                WA_LOOP(19);
+                WA_HALOAD_2(8);
+                WA_LOOP(20);
+                WA_HALOAD_2(9);
+                WA_LOOP(21);
+                WA_HALOAD_2(10);
+                WA_LOOP(22);
+                WA_HALOAD_2(11);
+                WA_LOOP(23);
+                WA_HALOAD_2(12);
+                WA_LOOP(24);
+                WA_HALOAD_2(13);
+                WA_LOOP(25);
+                WA_HALOAD_2(14);
+                WA_LOOP(26);
+                WA_HALOAD_2(15);
+
+                WA_LOOP(27);
+                WA_LOOP(28);
+                WA_LOOP(29);
+                WA_LOOP(30);
+                WA_LOOP(31);
+
+                __syncthreads();
+            }
+        }
+    }
+
+    if (scale) {
+        #pragma unroll
+        for (int c = 0; c < colorsPerThread; c++) {
+            #pragma unroll
+            for (int f = 0; f < filtersPerThread; f++) {
+                targets[c * B_Y * filterPixels * numFilters + f * B_X] = scaleTargets * targets[c * B_Y * filterPixels * numFilters + f * B_X] + scaleOutputs * prod[f][c];
+            }
+        }
+    } else {
+        #pragma unroll
+        for (int c = 0; c < colorsPerThread; c++) {
+            #pragma unroll
+            for (int f = 0; f < filtersPerThread; f++) {
+                targets[c * B_Y * filterPixels * numFilters + f * B_X] = scaleOutputs * prod[f][c];
+            }
+        }
+    }
+}
+
+/*
+ * images:      (numImgColors, imgSizeY, imgSizeX, numImages), with stride given
+ * hidActs:     (numFilters, numModulesY, numModulesX, numImages)
+ *
+ * targets:     (numModulesY*numModulesX/partialSum, numFilterColors, filterPixels, numFilters)
+ */
+template <int B_Y, int B_X, int filtersPerThread, int colorsPerThread, int preloadCases, bool scale>
+__launch_bounds__(256, 2)
+__global__ void conv_weight_acts_mc_mf_kepler_preload_ty_8_tx_32_f_4_c_8_r_16_tex(cudaTextureObject_t images, cudaTextureObject_t hidActs, float* targets,
                                        const int numImages, const int numFilters,
                                        const int numModulesY, const int numModulesX,
                                        const int imgSizeY, const int imgSizeX, const int filterSize,

            else if (numFilterColors <= 3) {
                if (numFilterColors == 3) {
                    if (numFiltersPerGroup % 64 == 0) {
                      cudaTextureObject_t texImages = THCudaTensor_getTextureObject(images);
                      cudaTextureObject_t texHidActs = THCudaTensor_getTextureObject(hidActs);
                        cudaFuncSetCacheConfig(conv_weight_acts_c_preload_pc_2_pt_2_f_4_r_32_c_3 < 16, 16, 2, 2, 4, 32, 3, false, false >, cudaFuncCachePreferShared);
                        conv_weight_acts_c_preload_pc_2_pt_2_f_4_r_32_c_3 < 16, 16, 2, 2, 4, 32, 3, false, false ><<<blocks, threads, 0>>>(texImages, texHidActs, THCudaTensor_data(targets), numImages, numFilters, numModulesY, numModulesX, imgSizeY, imgSizeX, filterSize, paddingStart, moduleStride, imgStride, sumWidth, scaleTargets, scaleOutput);
                      checkCudaErrors(cudaDestroyTextureObject(texImages));
                      checkCudaErrors(cudaDestroyTextureObject(texHidActs));
                    }
                    else if (numFiltersPerGroup % 48 == 0) {
                      cudaTextureObject_t texImages = THCudaTensor_getTextureObject(images);
                      cudaTextureObject_t texHidActs = THCudaTensor_getTextureObject(hidActs);
                        cudaFuncSetCacheConfig(conv_weight_acts_c_preload_pc_2_pt_4_f_3_r_32_c_3 < 16, 16, 2, 4, 3, 32, 3, false, false >, cudaFuncCachePreferShared);
                        conv_weight_acts_c_preload_pc_2_pt_4_f_3_r_32_c_3 < 16, 16, 2, 4, 3, 32, 3, false, false ><<<blocks, threads, 0>>>(texImages, texHidActs, THCudaTensor_data(targets), numImages, numFilters, numModulesY, numModulesX, imgSizeY, imgSizeX, filterSize, paddingStart, moduleStride, imgStride, sumWidth, scaleTargets, scaleOutput);
                      checkCudaErrors(cudaDestroyTextureObject(texImages));
                      checkCudaErrors(cudaDestroyTextureObject(texHidActs));
                    }
                    else if (numFiltersPerGroup % 32 == 0) {
                        cudaFuncSetCacheConfig(conv_weight_acts_c_kepler_sw < 8, 16, 2, 2, 2, 16, 3, false, false >, cudaFuncCachePreferShared);

 */

#include "../include/cudaconv2.cuh"
#include "../../nvmatrix/include/nvmatrix_kernels.cuh"

#define LO16(x)     ((x) & 0x0000FFFF)
#define HI16(x)     ((x) >> 16)
