Howdy, Stranger!

It looks like you're new here. If you want to get involved, click one of these buttons!

Categories

In this Discussion

Tags in this Discussion

trouble with cudaBindTexture2D/tex2D and toolkit 4.0/4.1
  • Vote Up0Vote Down franzdaubner
    Posts: 18
    I recently tried to switch all my kernels from toolkit 3.2 to toolkit 4.1.

    One of the kernels (the only one using textures) is not working anymore, it just behaves like an empty kernel with no code in it. There's no error message, no warning, it just does nothing. This behaviour is reproducible with toolkit 4.0

    If I switch back to toolkit 3.2 everything works fine.

    So here is my question:
    Did anyone use cudaBindTexture2D/tex2D with success on toolkit 4.1? On what hardware?
    I'm using a Geforce GTX285 on Windows 7 (64) and Visual Studio 2008.
  • 6 Comments sorted by
  • Vote Up1Vote Down fwende
    Posts: 19 Accepted Answer
    i have cuda 4.0 (270.41.19) installed on my workstation (2 x gtx285, ubuntu linux 64-bit). i used 2d-textures in one of my previous projects and it worked without any errors. device memory was allocated using cudaMalloc(), and then bound via cudaBindTexture2D().

    just now, i recompiled the code and checked simulation results, and it is still working. i also tested on a tesla m2090 (cuda 4.0, 270.41.19, scientific linux 64-bit).
  • Vote Up1Vote Down fwende
    Posts: 19
    did you try running a simple code sample that uses texture memory, or one of the sdk-samples that uses textures in a similar way you do (i don't know if there is any).
  • Vote Up1Vote Down franzdaubner
    Posts: 18 Accepted Answer
    fwende said:

    did you try running a simple code sample that uses texture memory, or one of the sdk-samples that uses textures in a similar way you do (i don't no if there is any).



    I didn't find any sample code that uses bindtexture2D, so I now wrote a kernel which just copies a memory block using tex2D.

    And I think I found out what causes the problems: I read an image in rgba layout (uchar4), and the format of the texture is float4 since I want to use linear interpolation. This works with toolkit 3.2, but not with 4.0 or 4.1

    To be more precise, this works:

    texture<uchar4, 2, cudaReadModeElementType> memCopyArea2D;
    ...
    cudaChannelFormatDesc chDescArea2D = cudaCreateChannelDesc<uchar4>();

    But this doesn't:

    texture<float4, 2, cudaReadModeElementType> memCopyArea2D;
    ...
    cudaChannelFormatDesc chDescArea2D = cudaCreateChannelDesc<uchar4>();
  • Vote Up1Vote Down franzdaubner
    Posts: 18 Accepted Answer
    I had a very enlightening email-conversation with a support-technician from nvidia, and things have been cleared up.

    When I wrote the kernel which caused the problems, I found that the documentation of the texture-units is very unclear at best. So I did some trial and error until I found a working combination, and drew my conclusions on how texture units work based on that. Turns out that there was actually a bug in the old toolkit which allowed illegal combinations, so my conclusions were wrong.

    I thought that by declaring the texture-datatype as float4 and the channel description as uchar4 I could get the texture-unit to perform an int-to-float conversion (since I wanted to use linear interpolation on the uchar4-values).

    But the only way to do that is to declare the read-mode as cudaReadModeNormalizedFloat.
    Both the declaration of the texture-datatype and the declaration of the channel descriptor have to be the same, since they both refer to the memory layout of the raw data. This design is of course redundant and confusing, even more so since there is no check if the data type of the texture and the channel descriptor match.

    So this is the correct way to get the texture-unit to use bilinear interpolation on byte-values (32-bit image or uchar4):

    texture<uchar4, 2, cudaReadModeNormalizedFloat> memCopyArea2D;
    ...
    memCopyArea2D.filterMode = cudaFilterModeLinear;
    ...
    cudaChannelFormatDesc chDescArea2D = cudaCreateChannelDesc<uchar4>();

    A call to tex2D will then return the type float4, and the bilinear interpolation will work as expected.
  • Vote Up0Vote Down franzdaubner
    Posts: 18
    I already filed a bug for this (#74516), and it has been reproduced by the nvidia team. I would still be interested if anyone uses cudaBindTexture2D with toolkit 4.0 / 4.1 with success.
    Or am I the only one using textures without cudaArrays? :)
  • Vote Up0Vote Down franzdaubner
    Posts: 18
    fwende said:

    just now, i recompiled the code and checked simulation results, and it is still working.



    Hmm, so maybe it's something else which is causing the problems... or it's windows specific. Oh well, let's see what the nvidia team comes up with :)

    Thank you for the feedback...
Add a Comment