is Texture memory Read-only? is texture memory in fact read-only or not?

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi folk,

although it is stated throughout all manuals and forum topics that the texture memory is read-only, i still have my doubts. The familiar example on the texture memory usage from the “CUDA by example” (chapter 7, heat transfer) book made me feel that in fact there was writing into the texture memory. So i have modified (or rather simplified) that example and, as you can see, after 3 calls to the ‘copy_const_kernel’ function the content of the texture memory is different from what it initially was. So, is in fact texture memory is NOT read-only or i have misinterpreted something? Thanks for any suggestions.

program code

#include <stdio.h>
#define N 10

texture tex;

int dataSize = N*sizeof(float);
float res = (float)malloc( dataSize );

global void copy_const_kernel( float *iptr ) {

int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
int offset = x + y * blockDim.x * gridDim.x;
if (offset<N)
   iptr[offset] = tex1Dfetch(tex,offset)+0.1;

}

int main( void ) {

float *devtxt;

cudaMalloc( (void**)&devtxt, dataSize );
cudaBindTexture( NULL, tex, devtxt, dataSize );

// intialize the constant data
float *temp = (float*)malloc( dataSize );
for (int i=0; i<N; i++) temp[i] = (float)i;

cudaMemcpy( devtxt, temp, dataSize, cudaMemcpyHostToDevice );    

free( temp );

cudaMemcpy( res, devtxt, dataSize, cudaMemcpyDeviceToHost );

for (int i=0; i<N; i++) printf("before: %f \n",res[i]);

copy_const_kernel<<<16,4>>>( devtxt );
copy_const_kernel<<<16,4>>>( devtxt );
copy_const_kernel<<<16,4>>>( devtxt );

cudaMemcpy( res, devtxt, dataSize, cudaMemcpyDeviceToHost );

for (int i=0; i<N; i++) printf("after: %f \n",res[i]);

cudaUnbindTexture( tex );
cudaFree( devtxt );
free(res);

}

program output:

before: 0.000000
before: 1.000000
before: 2.000000
before: 3.000000
before: 4.000000
before: 5.000000
before: 6.000000
before: 7.000000
before: 8.000000
before: 9.000000
after: 0.300000
after: 1.300000
after: 2.300000
after: 3.300000
after: 4.300000
after: 5.300000
after: 6.300000
after: 7.300000
after: 8.300001
after: 9.300001

2

hi!

you are writing to the global memory which is bound to texture reference (not to the texture directly).
Texture is a special mechanism to access global memory. Texture has only one type of memory - cache,
and you can’t write something into that cache (e.g. tex1D(ref, x) = 12.0 is illegal).

3

I think from reading the docs (I’ve never tried it myself) that you can write to the underlying array, but because texture references are cached, the changes aren’t guaranteed to show up until you run the next kernel. But in your example, you’ve got three kernel invocations, and the texture cache is probably flushed on each invocation, so it works. Try a read-write-read within a kernel…