cudaBindTexture2D showing error.

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,
I am working on GeForce GT 610.
When I tried to bind memory to texture refernece , I am getting error.
I am pasting code below.
I am getting error on line no 37.

#include <opencv2/core/core.hpp>
#include <opencv2/highgui/highgui.hpp>
#include <iostream>
#include "cuda.h"
#include <stdio.h>

using namespace cv;
using namespace std;

#define checkCudaErrors(err)  __checkCudaErrors (err, __FILE__, __LINE__)

inline void __checkCudaErrors(cudaError err, const char *file, const int line )
{
    if(cudaSuccess != err)
    {
        fprintf(stderr, "%s(%i) : CUDA Runtime API error %d: %s.\n",file, line, (int)err, cudaGetErrorString( err ) );
        exit(-1);
    }
}

texture<float,2> myTex;

int main(int argc, char* argv[])
{
    float* input;
    input = new float[348 * 48];
    for(int i = 0; i < 348 * 48; ++i)
    {
        input[i] = i;
    }
    float* inputDevice;
    checkCudaErrors(cudaMalloc ((void**)&inputDevice, 348 * 48 * sizeof(float) ));
    checkCudaErrors(cudaMemcpy(inputDevice, input, 348 * 48 * sizeof(float), cudaMemcpyHostToDevice));

    cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();

    checkCudaErrors(cudaBindTexture2D(0, myTex, inputDevice, desc, 348, 48, sizeof(float) * 348));

    cudaUnbindTexture(myTex);
    cudaFree(inputDevice);
    cout << "done" << endl;

    return 0;
}

Thanks,
Tushar

2

Can anyone please help me understand why this error is coming?
Error message: invalid argument

Thanks,
Tushar

3

cudaBindTexture2D(0, myTex, inputDevice, desc, 348, 48, sizeof(float) * 348)

cudaError_t cudaBindTexture2D (size_t *offset, const
textureReference *texref, const void *devPtr, const
cudaChannelFormatDesc *desc, size_t width, size_t
height, size_t pitch)

cudaChannelFormatDesc is probably the only parameter passed correctly

4

Generally speaking, 2D texturing requires a properly pitched allocation.

The following modifications to your code represent one possible form of valid usage:

#include <iostream>
#include <stdio.h>

using namespace std;

#define checkCudaErrors(err)  __checkCudaErrors (err, __FILE__, __LINE__)

inline void __checkCudaErrors(cudaError err, const char *file, const int line )
{
    if(cudaSuccess != err)
    {
        fprintf(stderr, "%s(%i) : CUDA Runtime API error %d: %s.\n",file, line, (int)err, cudaGetErrorString( err ) );
        exit(-1);
    }
}

texture<float,2> myTex;

int main(int argc, char* argv[])
{
    float* input;
    input = new float[348 * 48];
    for(int i = 0; i < 348 * 48; ++i)
    {
        input[i] = i;
    }
    float* inputDevice;
    size_t pitch;
    checkCudaErrors(cudaMallocPitch((void**)&inputDevice, &pitch, 348 * sizeof(float),48 ));
    checkCudaErrors(cudaMemcpy2D(inputDevice, pitch, input, 348 * sizeof(float),348*sizeof(float), 48,  cudaMemcpyHostToDevice));

    cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();

    checkCudaErrors(cudaBindTexture2D(0, myTex, inputDevice, desc, 348, 48, pitch));

    cudaUnbindTexture(myTex);
    cudaFree(inputDevice);
    cout << "done" << endl;

    return 0;
}

You may also wish to refer to the cuda pitch linear texture sample code:

http://docs.nvidia.com/cuda/cuda-samples/index.html#pitch-linear-texture

which also gives an example of proper usage of cudaBindTexture2D

5

Hi txbob,

Thanks for your reply.
I tried to run your edited code, it is giving me value of pitch = 384 ( 128 *3).
Its taking pitch values in terms of 128 for memory coalescing.

When in my code when I tried to make following changes in line no 37

checkCudaErrors(cudaBindTexture2D(0, myTex, inputDevice, desc, 348, 48, sizeof(float) * 348));
To
checkCudaErrors(cudaBindTexture2D(0, myTex, inputDevice, desc, 348, 48, sizeof(float) * 352));

Now its not showing any error message. Does it mean that pitch value may be in multiples of 8 ?(352 is multiple of 8).I tried same with numbers which are multiples of 8.

Thanks,
Tushar

6

pitch is not intended to be something you set yourself. pitch is intended to be returned by the runtime, and used as-is. It’s remotely possible that pitch mechanics could change on future hardware.

The single line you have changed would be inconsistent, pitch-wise, with the previous allocation, most likely. The runtime may not be detecting an error at this point, but I’m not sure the texturing results would be correct in all cases.

Probably it is beyond my expertise to give you a full answer. I would recommend using the pitch parameter the way it was intended to be used, AFAIK.