     dim3 blocks2 = dim3(DIVUP(numImages,32*4) * imgSize, (numFilters / 4) * imgSize);
 
     bool checkCaseBounds = (numImages % 128) != 0;
+    cudaTextureObject_t texOutGrads = THCudaTensor_getTextureObject(outGrads);
+    cudaTextureObject_t texInputs = THCudaTensor_getTextureObject(inputs);
+    cudaTextureObject_t texActs = THCudaTensor_getTextureObject(acts);
     if (blocked) {
         if (scaleTargets == 0 && scaleOutput == 1) {
             if (checkCaseBounds) {
                 cudaFuncSetCacheConfig(kFRNormUndo2<4, 32, 4, false, true, true>, cudaFuncCachePreferL1);
-                kFRNormUndo2<4, 32, 4, false, true, true><<<blocks2, threads2, 0>>>(THCudaTensor_getTextureObject(outGrads), THCudaTensor_getTextureObject(inputs), THCudaTensor_getTextureObject(acts),
+                kFRNormUndo2<4, 32, 4, false, true, true><<<blocks2, threads2, 0>>>(texOutGrads, texInputs, texActs,
                                                                         THCudaTensor_data(target), imgSize, numFilters, numImages, sizeF, addScale, powScale, minDiv,
                                                                         scaleTargets, scaleOutput);
             } else {
                 cudaFuncSetCacheConfig(kFRNormUndo2<4, 32, 4, false, false, true>, cudaFuncCachePreferL1);
-                kFRNormUndo2<4, 32, 4, false, false, true><<<blocks2, threads2, 0>>>(THCudaTensor_getTextureObject(outGrads), THCudaTensor_getTextureObject(inputs), THCudaTensor_getTextureObject(acts),
+                kFRNormUndo2<4, 32, 4, false, false, true><<<blocks2, threads2, 0>>>(texOutGrads, texInputs, texActs,
                                                                         THCudaTensor_data(target), imgSize, numFilters, numImages, sizeF, addScale, powScale, minDiv,
                                                                         scaleTargets, scaleOutput);
             }
         } else {
             if (checkCaseBounds) {
                 cudaFuncSetCacheConfig(kFRNormUndo2<4, 32, 4, true, true, true>, cudaFuncCachePreferL1);
-                kFRNormUndo2<4, 32, 4, true, true, true><<<blocks2, threads2, 0>>>(THCudaTensor_getTextureObject(outGrads), THCudaTensor_getTextureObject(inputs), THCudaTensor_getTextureObject(acts),
+                kFRNormUndo2<4, 32, 4, true, true, true><<<blocks2, threads2, 0>>>(texOutGrads, texInputs, texActs,
                                                                         THCudaTensor_data(target), imgSize, numFilters, numImages, sizeF, addScale, powScale, minDiv,
                                                                         scaleTargets, scaleOutput);
             } else {
                 cudaFuncSetCacheConfig(kFRNormUndo2<4, 32, 4, true, false, true>, cudaFuncCachePreferL1);
-                kFRNormUndo2<4, 32, 4, true, false, true><<<blocks2, threads2, 0>>>(THCudaTensor_getTextureObject(outGrads), THCudaTensor_getTextureObject(inputs), THCudaTensor_getTextureObject(acts),
+                kFRNormUndo2<4, 32, 4, true, false, true><<<blocks2, threads2, 0>>>(texOutGrads, texInputs, texActs,
                                                                         THCudaTensor_data(target), imgSize, numFilters, numImages, sizeF, addScale, powScale, minDiv,
                                                                         scaleTargets, scaleOutput);
             }

        if (scaleTargets == 0 && scaleOutput == 1) {
            if (checkCaseBounds) {
                cudaFuncSetCacheConfig(kFRNormUndo2<4, 32, 4, false, true, false>, cudaFuncCachePreferL1);
                kFRNormUndo2<4, 32, 4, false, true, false><<<blocks2, threads2, 0>>>(THCudaTensor_getTextureObject(outGrads), THCudaTensor_getTextureObject(inputs), THCudaTensor_getTextureObject(acts),
                                                                        THCudaTensor_data(target), imgSize, numFilters, numImages, sizeF, addScale, powScale, minDiv,
                                                                        scaleTargets, scaleOutput);
            } else {
                cudaFuncSetCacheConfig(kFRNormUndo2<4, 32, 4, false, false, false>, cudaFuncCachePreferL1);
                kFRNormUndo2<4, 32, 4, false, false, false><<<blocks2, threads2, 0>>>(THCudaTensor_getTextureObject(outGrads), THCudaTensor_getTextureObject(inputs), THCudaTensor_getTextureObject(acts),
                                                                        THCudaTensor_data(target), imgSize, numFilters, numImages, sizeF, addScale, powScale, minDiv,
                                                                        scaleTargets, scaleOutput);
            }
        } else {
            if (checkCaseBounds) {
                cudaFuncSetCacheConfig(kFRNormUndo2<4, 32, 4, true, true, false>, cudaFuncCachePreferL1);
                kFRNormUndo2<4, 32, 4, true, true, false><<<blocks2, threads2, 0>>>(THCudaTensor_getTextureObject(outGrads), THCudaTensor_getTextureObject(inputs), THCudaTensor_getTextureObject(acts),
                                                                        THCudaTensor_data(target), imgSize, numFilters, numImages, sizeF, addScale, powScale, minDiv,
                                                                        scaleTargets, scaleOutput);
            } else {
                cudaFuncSetCacheConfig(kFRNormUndo2<4, 32, 4, true, false, false>, cudaFuncCachePreferL1);
                kFRNormUndo2<4, 32, 4, true, false, false><<<blocks2, threads2, 0>>>(THCudaTensor_getTextureObject(outGrads), THCudaTensor_getTextureObject(inputs), THCudaTensor_getTextureObject(acts),
                                                                        THCudaTensor_data(target), imgSize, numFilters, numImages, sizeF, addScale, powScale, minDiv,
                                                                        scaleTargets, scaleOutput);
            }
        }
    }

    getLastCudaError("convResponseNormCrossMapUndo: kernel execution failed");
}


    dim3 blocks(DIVUP(THCudaTensor_nElement(outGrads),(threads.x * prelimEltsPerThread)));
    bool checkPrelimBounds = THCudaTensor_nElement(outGrads) % (threads.x * prelimEltsPerThread) != 0;
    //printf("num elts: %d, blocks: %d\n", outGrads.getNumElements(), blocks.x);
    cudaTextureObject_t texDenoms = THCudaTensor_getTextureObject(denoms);
    cudaTextureObject_t texOutGrads = THCudaTensor_getTextureObject(outGrads);
    kRNormUndoPrelims<128, 8><<<blocks, threads, 0>>>(THCudaTensor_data(acts), texDenoms, texOutGrads, THCudaTensor_nElement(outGrads), -2*addScale*powScale);
    checkCudaErrors(cudaDestroyTextureObject(texDenoms));
    checkCudaErrors(cudaDestroyTextureObject(texOutGrads));

    // Now the main routine
    if (sizeX >= 6 && numFilters % 4 == 0) {
