Cuda understanding texture fetching

Question

I have written a simple code example i have cuda version 6.5

#include <iostream>
#include <cstdio>
#include "cudaerror.h"
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

texture<unsigned int, cudaTextureType2D, cudaReadModeElementType> texRef;

__global__
void kernel1(int64_t N){
    int64_t tid_x = threadIdx.x;
    int64_t tid_y = threadIdx.y;

    if(tid_x < N && tid_y < N){
        unsigned int temp = tex2D(texRef, tid_x, tid_y);
        printf("tid_x: %d, tid_y: %d, tex: %d\n", tid_x, tid_y, temp);
    }
}

void alloc_darrays(cudaArray* &d_adj_mat){
    unsigned int* adj_mat = new unsigned int[9]();

    adj_mat[3] = 1;
    adj_mat[4] = 1;
    adj_mat[5] = 1;
    adj_mat[0] = 1;

    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32,0,0,0,cudaChannelFormatKindUnsigned);

    gpuErrchk( cudaMallocArray(&d_adj_mat, &channelDesc, 3, 3) );
    gpuErrchk( cudaMemcpyToArray(d_adj_mat, 0, 0, adj_mat, 9*sizeof(int), cudaMemcpyHostToDevice) );

    texRef.addressMode[0] = cudaAddressModeBorder;
    texRef.addressMode[1] = cudaAddressModeBorder;
    texRef.filterMode = cudaFilterModePoint;
    texRef.normalized = false;

    gpuErrchk( cudaBindTextureToArray(texRef, d_adj_mat, channelDesc) );
}

int main(){
    cudaArray* d_adj_mat;
    alloc_darrays(d_adj_mat);
    dim3 numthreads(3,3);
    kernel1<<<1,numthreads>>>(3);
    cudaDeviceSynchronize();
    return 0;
}

The output produced is :

tid_x: 0, tid_y: 0, tex: 0
tid_x: 1, tid_y: 0, tex: 0
tid_x: 2, tid_y: 0, tex: 0
tid_x: 0, tid_y: 0, tex: 1
tid_x: 1, tid_y: 0, tex: 1
tid_x: 2, tid_y: 0, tex: 1
tid_x: 0, tid_y: 0, tex: 2
tid_x: 1, tid_y: 0, tex: 2
tid_x: 2, tid_y: 0, tex: 2

I am unable to understand from where the 2 is coming from and why tid_y always is 0.

Is there somthing wrong with the code or am i understanding textures wrong?

No, i didn't try them out yet, but i was just doing this and noticed this behaviour and couldn't understand it — Artemis Fowl
– Artemis Fowl, Commented Oct 23, 2015 at 10:35
your code is not a minimal reproducible example, since cudaerror.h is missing. You should provide code that can be copied, pasted and compiled without any modifications (ideally in one single, self contained file) — m.s.
– m.s., Commented Oct 23, 2015 at 10:36
Sorry it just contained a macro for error handling gpuErrchk i will remove it. — Artemis Fowl
– Artemis Fowl, Commented Oct 23, 2015 at 10:56

2 revs, 2 users 93% · Accepted Answer · 2015-10-23 10:52:34Z

1

What you are seeing is an artifact of using an incorrect kernel printf format string.

What you probably want to do is this:

printf("tid_x: %ld, tid_y: %ld, tex: %u\n", tid_x, tid_y, temp);

because your tid_x and tid_y variables are long integers. If you change that, the kernel should work as expected.

edited Oct 23, 2015 at 10:52

community wiki

2 revs, 2 users 93%
talonmies

Sign up to request clarification or add additional context in comments.

1 Comment

Artemis Fowl Over a year ago

yeah pretty stupid of me not to notice, probably because i didn't get compiler warnings

Collectives™ on Stack Overflow

Cuda understanding texture fetching

1 Answer 1

1 Comment

Your Answer

Hot Network Questions

Collectives™ on Stack Overflow

1 Answer 1

1 Comment

Your Answer

Sign up or log in

Post as a guest

Related