             else if (numFilterColors <= 3) {
                 if (numFilterColors == 3) {
                     if (numFiltersPerGroup % 64 == 0) {
-                      cudaTextureObject_t texImages = THCudaTensor_getTextureObject(images);
-                      cudaTextureObject_t texHidActs = THCudaTensor_getTextureObject(hidActs);
+                      if (THCudaTensor_nElement(images) * 4 < TEXTURE_SIZE_MAX) {
+                        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_tex < 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_tex < 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 {
                         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));
+                        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>>>(THCudaTensor_data(images), THCudaTensor_data(hidActs), THCudaTensor_data(targets), numImages, numFilters, numModulesY, numModulesX, imgSizeY, imgSizeX, filterSize, paddingStart, moduleStride, imgStride, sumWidth, scaleTargets, scaleOutput);
+                      }
                     }
                     else if (numFiltersPerGroup % 48 == 0) {
-                      cudaTextureObject_t texImages = THCudaTensor_getTextureObject(images);
-                      cudaTextureObject_t texHidActs = THCudaTensor_getTextureObject(hidActs);
+                      if (THCudaTensor_nElement(images) * 4 < TEXTURE_SIZE_MAX) {
+                        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_tex < 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_tex < 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 {
                         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));
+                        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>>>(THCudaTensor_data(images), THCudaTensor_data(hidActs), THCudaTensor_data(targets), numImages, numFilters, numModulesY, numModulesX, imgSizeY, imgSizeX, filterSize, paddingStart, moduleStride, imgStride, sumWidth, scaleTargets, scaleOutput);
+                      }
                     }
                     else if (numFiltersPerGroup % 32 == 0) {
                         cudaFuncSetCacheConfig(conv_weight_acts_c_kepler_sw < 8, 16, 2, 2, 2, 16, 3, false, false >, cudaFuncCachePreferShared);

 * 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_6_r_32(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,

 */
template <int B_Y, int B_X, int pixelCache, int pixelsPerThread, int filtersPerThread, int preloadCases, int numColors, bool scale, bool checkCaseBounds>
//__launch_bounds__(256,2)
__global__ void conv_weight_acts_c_preload_pc_2_pt_2_f_4_r_32_c_3_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,
