     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);
-    kRNormUndoPrelims<128, 8><<<blocks, threads, 0>>>(THCudaTensor_data(acts), THCudaTensor_getTextureObject(denoms), THCudaTensor_getTextureObject(outGrads), THCudaTensor_nElement(outGrads), -2*addScale*powScale);
+    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) {

        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(numImages,32*4) * imgSize, (numFilters / 4) * imgSize);
//    printf("convContrastNormCrossMap imgs: %p, meanDiffs: %p, denoms: %p, target: %p, imgSize: %d, numFilters: %d, numImages: %d, sizeF: %d, addScale: %f, powScale: %f, minDiv: %f, blocked: %d\n",
//            THCudaTensor_data(images), THCudaTensor_data(meanDiffs), THCudaTensor_data(denoms), THCudaTensor_data(target), imgSize, numFilters, numImages, sizeF, addScale, powScale, minDiv, blocked);
    cudaTextureObject_t texImages = THCudaTensor_getTextureObject(images);
    cudaTextureObject_t texMeanDiffs = THCudaTensor_getTextureObject(meanDiffs);
    if (blocked) {
        if (checkCaseBounds) {
            cudaFuncSetCacheConfig(kFCNorm<4, 32, 4, true, true>, cudaFuncCachePreferL1);
            kFCNorm<4, 32, 4, true, true><<<blocks, threads, 0>>>(texImages, texMeanDiffs, THCudaTensor_data(target),
                                                                imgSize, numFilters, numImages, sizeF, addScale, powScale, minDiv);
        } else {
            cudaFuncSetCacheConfig(kFCNorm<4, 32, 4, false, true>, cudaFuncCachePreferL1);
            kFCNorm<4, 32, 4, false, true><<<blocks, threads, 0>>>(texImages, texMeanDiffs, THCudaTensor_data(target),
                                                                imgSize, numFilters, numImages, sizeF, addScale, powScale, minDiv);
        }
    } else {
        if (checkCaseBounds) {
            cudaFuncSetCacheConfig(kFCNorm<4, 32, 4, true, false>, cudaFuncCachePreferL1);
            kFCNorm<4, 32, 4, true, false><<<blocks, threads, 0>>>(texImages, texMeanDiffs, THCudaTensor_data(target),
                                                                imgSize, numFilters, numImages, sizeF, addScale, powScale, minDiv);
        } else {
            cudaFuncSetCacheConfig(kFCNorm<4, 32, 4, false, false>, cudaFuncCachePreferL1);
            kFCNorm<4, 32, 4, false, false><<<blocks, threads, 0>>>(texImages, texMeanDiffs, THCudaTensor_data(target),
                                                                imgSize, numFilters, numImages, sizeF, addScale, powScale, minDiv);
        }
    }
    checkCudaErrors(cudaDestroyTextureObject(texImages));
    checkCudaErrors(cudaDestroyTextureObject(texMeanDiffs));

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