Texture fetch slower than direct global access, chapter 7 from "Cuda by example" book

Question

I am reading and testing the examples in the book "Cuda By example. An introduction to General Purpose GPU Programming". When testing the examples in chapter 7, relative to texture memory, I realized that access to global memory via texture cache is much slower than direct access (My NVIDIA GPU is GeForceGTX 260, compute capability 1.3 and I am using NVDIA CUDA 4.2):

• Time per frame with texture fetch (1D or 2D) for a 256*256 image: 93 ms
• Time per frame not using texture (just direct global access) for 256*256: 8.5 ms

I have double checked the code several times and I have also been reading the "CUDA C Programming guide" and "CUDA C Best practices Guide" which come along with the SDK, and I do not really understand the problem. As far as I understand, texture memory is just global memory with a specific access mechanism implementation to make it look like a cache (?). I am wondering whether coalesced access to global memory will make texture fetch slower, but I cannot be sure.

Does anybody have a similar problem? (I found some links in NVIDIA forums for a similar problem, but the link is no longer available.)

The testing code looks this way, only including the relevant parts:

//#define TEXTURE
//#define TEXTURE2

#ifdef TEXTURE
// According to C programming guide, it should be static (3.2.10.1.1)
static texture<float> texConstSrc;
static texture<float> texIn;
static texture<float> texOut;
#endif

    __global__ void copy_const_kernel( float *iptr
    #ifdef TEXTURE2
     ){
    #else
        ,const float *cptr ) {
    #endif
            // map from threadIdx/BlockIdx to pixel position
            int x = threadIdx.x + blockIdx.x * blockDim.x;
            int y = threadIdx.y + blockIdx.y * blockDim.y;
            int offset = x + y * blockDim.x * gridDim.x;

    #ifdef TEXTURE2
            float c = tex1Dfetch(texConstSrc,offset);
    #else
            float c = cptr[offset];     
    #endif

            if ( c != 0) iptr[offset] = c;
    }

    __global__ void blend_kernel( float *outSrc,
    #ifdef TEXTURE
        bool dstOut ) {
    #else
        const float *inSrc ) {
    #endif
            // map from threadIdx/BlockIdx to pixel position
            int x = threadIdx.x + blockIdx.x * blockDim.x;
            int y = threadIdx.y + blockIdx.y * blockDim.y;
            int offset = x + y * blockDim.x * gridDim.x;
            int left = offset - 1;
            int right = offset + 1;
            if (x == 0) left++;
            if (x == SXRES-1) right--;
            int top = offset - SYRES;
            int bottom = offset + SYRES;
            if (y == 0) top += SYRES;
            if (y == SYRES-1) bottom -= SYRES;

    #ifdef TEXTURE
            float t, l, c, r, b;
            if (dstOut) {
                t = tex1Dfetch(texIn,top);
                l = tex1Dfetch(texIn,left);
                c = tex1Dfetch(texIn,offset);
                r = tex1Dfetch(texIn,right);
                b = tex1Dfetch(texIn,bottom);
            } else {
                t = tex1Dfetch(texOut,top);
                l = tex1Dfetch(texOut,left);
                c = tex1Dfetch(texOut,offset);
                r = tex1Dfetch(texOut,right);
                b = tex1Dfetch(texOut,bottom);
            }
            outSrc[offset] = c + SPEED * (t + b + r + l - 4 * c);
    #else
            outSrc[offset] = inSrc[offset] + SPEED * ( inSrc[top] +
                inSrc[bottom] + inSrc[left] + inSrc[right] -
                inSrc[offset]*4);
    #endif
    }

    // globals needed by the update routine
    struct DataBlock {
        unsigned char *output_bitmap;
        float *dev_inSrc;
        float *dev_outSrc;
        float *dev_constSrc;
        cudaEvent_t start, stop;
        float totalTime;
        float frames;
        unsigned size;
        unsigned char *output_host;
    };
    void anim_gpu( DataBlock *d, int ticks ) {
        checkCudaErrors( cudaEventRecord( d->start, 0 ) );
        dim3 blocks(SXRES/16,SYRES/16);
        dim3 threads(16,16);

    #ifdef TEXTURE
        volatile bool dstOut = true;
    #endif

        for (int i=0; i<90; i++) {
    #ifdef TEXTURE
            float *in, *out;
            if (dstOut) {
                in = d->dev_inSrc;
                out = d->dev_outSrc;
            } else {
                out = d->dev_inSrc;
                in = d->dev_outSrc;
            }
    #ifdef TEXTURE2
            copy_const_kernel<<<blocks,threads>>>( in );
    #else
            copy_const_kernel<<<blocks,threads>>>( in,
                d->dev_constSrc );
    #endif
            blend_kernel<<<blocks,threads>>>( out, dstOut );
            dstOut = !dstOut;

    #else
            copy_const_kernel<<<blocks,threads>>>( d->dev_inSrc,
                d->dev_constSrc );
            blend_kernel<<<blocks,threads>>>( d->dev_outSrc,
                d->dev_inSrc );
            swap( d->dev_inSrc, d->dev_outSrc );
    #endif
        }
            // Some stuff for the events
            // ...
         }

Answer 1

I have been testing the results with the nvvp (NVIDIA profiler)

The result are quite curious as they show that there are a lot of texture cache misses (which are probably the cause for the bad performance). The result from the profiler show also information that is difficult to understand even using the guide "CUPTI_User_GUide):

• text_cache_hit: Number of texture cache hits (they are accounted only for one SM according to 1.3 capability).

• text_cache_miss: Number of texture cache miss (they are accounted only for one SM according to 1.3 capability).

The following are the results for an example of 256*256 without using texture cache (only relevant info is shown):

Name Duration(ns) Grid_Size Block_Size

"copy_const_kernel(...) 22688 16,16,1 16,16,1

"blend_kernel(...)" 51360 16,16,1 16,16,1

Following are the results using 1D texture cache:

Name Duration(ns) Grid_Size Block_Size tex_cache_hit tex_cache_miss

"copy_const_kernel(...)" 147392 16,16,1 16,16,1 0 1024

"blend_kernel(...)" 841728 16,16,1 16,16,1 79 5041

Following are the results using 2D texture cache:

Name Duration(ns) Grid_Size Block_Size tex_cache_hit tex_cache_miss

"copy_const_kernel(...)" 150880 16,16,1 16,16,1 0 1024

"blend_kernel(...)" 872832 16,16,1 16,16,1 2971 2149

These result show several interesting info:

• There are no cache hits at all for the "copy const" function (although ideally the memory is "spatially located", in the sense that each thread accesses memory which is near to the memory acceded by other near threads). I guess that this is because the threads within this function do not access memory from other threads, which seems to be the way for the texture cache to be usable (being the "spatially located" concept quite confusing)

• There are some cache hits in the 1D and a lot more in the 2D case for the function "blend_kernel". I guess that it is due to the fact that within that function, any thread access memory from their neighbours threads. I cannot understand why there are more in 2D than 1d.

• The duration time is greater in the texture cases than in the no texture case (nearly about one order of magnitude). Perhaps related with the so many texture cache misses.

• For the "copy_const" function there are 1024 total accesses for the SM and 5120 for the "blend kernel". The relation 5:1 is correct due to the fact that there are 5 fetches in "blend" and only 1 in "copy_const". Anyway, I cannot understand where all this 1024 come from: ideally, this event "text cache miss/hot" only accounts for one SM (I have 24 in my GeForceGTX 260) and it only accounts for warps ( 32 thread size). Therefore, I have 256 threads/32=8 warps per SM and 256 blocks/24 = 10 or 11 "iterations" per SM, so I would be expecting something like 80 or 88 fetches (more over, some other event like sm_cta_launched, which is the number of thread blocks per SM, which is supposed to be supported in my 1.3 device, is always 0...)