         }
 
         __syncthreads();
-
-        FA_COLOR3_IMPRELOAD(0,0);
-        FA_COLOR3_IMPRELOAD(0,1);
-        FA_COLOR3_IMPRELOAD(0,2);
-        FA_COLOR3_IMPRELOAD(0,3);
-
+        FA_COLOR3_IMPRELOAD_TX(1,0);
+        FA_COLOR3_IMPRELOAD_TX(1,1);
+        FA_COLOR3_IMPRELOAD_TX(1,2);
+        FA_COLOR3_IMPRELOAD_TX(1,3);
+        FA_COLOR3_IMPRELOAD_TX(2,0);
+        FA_COLOR3_IMPRELOAD_TX(2,1);
+        FA_COLOR3_IMPRELOAD_TX(2,2);
+        FA_COLOR3_IMPRELOAD_TX(2,3);
         #pragma unroll
         for (int c = 0; c < numColors; c++) {
             #pragma unroll
             for (int pp = 0; pp < 2; pp++) {
+                fPreload[c][pp] = fPidxNext + pp*(B_X/filtersPerThread) + shFilterLoadY >= filterPixels ? 0 : tex1Dfetch<float>(filters, filtersOffset2 + c * numFilters* filterPixels + pp*(B_X/filtersPerThread) * numFilters);
+            }
+        }
+        #pragma unroll
+        for (int pp = 0; pp < pixelCache; pp++) {
+            #pragma unroll
+            for (int c = 0; c < numColors; c++) {
                 #pragma unroll
-                for(int i = 0; i < imgsPerThread; i++) {
+                for(int f = 0; f < filtersPerThread; f++) {
                     #pragma unroll
-                    for(int f = 0; f < filtersPerThread; f++) {
+                    for(int i = 0; i < imgsPerThread; i++) {
                         prod[i][f] += shImages[c][pp][tx * imgsPerThread + i] * shFilters[c][pp][ty * filtersPerThread + f];
                     }
                 }
-            }
-        }
-
-        FA_COLOR3_IMPRELOAD(2,0);
-        FA_COLOR3_IMPRELOAD(2,1);
-        FA_COLOR3_IMPRELOAD(2,2);
-        FA_COLOR3_IMPRELOAD(2,3);
-
-        #pragma unroll
-        for (int c = 0; c < numColors; c++) {
-            #pragma unroll
-            for (int pp = 2; pp < pixelCache; pp++) {
-                #pragma unroll
-                for(int i = 0; i < imgsPerThread; i++) {
-                    #pragma unroll
-                    for(int f = 0; f < filtersPerThread; f++) {
-                        prod[i][f] += shImages[c][pp][tx * imgsPerThread + i] * shFilters[c][pp][ty * filtersPerThread + f];
-                    }
-                }
-            }
-            #pragma unroll
-            for (int pp = 0; pp < 2; pp++) {
-                fPreload[c][pp] = fPidxNext + pp*(B_X/filtersPerThread) + shFilterLoadY >= filterPixels ? 0 : ff[c * numFilters* filterPixels + pp*(B_X/filtersPerThread) * numFilters];
             }
         }
 

        #pragma unroll
        for (int p = 0; p < pixelCache; p += B_X/filtersPerThread) {
            if (p + shFilterLoadY < filterPixels) {
                fPreload[c][p*filtersPerThread/B_X] = filters[p * numFilters + c * numFilters * filterPixels];
            } else{
                fPreload[c][p*filtersPerThread/B_X] = 0;
            }


    for (int p = 0; p < filterPixels; p += pixelCache) {
        #pragma unroll
        for (int i = 0; i < imgsPerThread; i++) {
            #pragma unroll
            for (int c = 0; c < numColors; ++c) {
                // NOTE: bank conflicts here!
                shImages[c][ty][tx * imgsPerThread + i] = imPreload[c][i];
            }
