Tex3D Optix 7

Gaming and Visualization Technologies Visualization OptiX
1

Good Morning,

I have some Optix 5/6 code that employs textures in a shader via CUDA tex3D call. The following is a snippet of older code I am trying to convert:

rtTextureSampler<float4, 3> myTexture;
RT_PROGRAM void closestHit()
{
// get data for T (float3)
float4 s1;
s1 = tex3D(myTexture,T.x,T.y,T.z);
}

So far I have converted to OptiX 7 as in following where the first part is in a separate header file called system.h
struct Params { float4 *myTexture; };

Second part is file called myOptixTexture.cu:

include <optix.h>
include “system.h”
extern "C" __constant__ Params spm;
extern "C" __global__ void __closesthit__ch() {

// get data for T (float3)
float4 s1;
s1 = tex3D(spm.myTexture,T.x,T.y,T.z);
}

The above s1 = tex3D(spm.myTexture,T.x,T.y,T.z); complains with the following:

no instance of overloaded function “tex3D” matches the argument list
argument types are: (float4 *, float, float, float)

Do I need to define the spm.myTexture differently? Is there a more standard way to approach retrieving texture(s) in OptiX 7?

Thanks for any help or hints and I apologize for the code formatting issues I have with posting to this forum - I can’t seem to get multiple line code to fall in the ‘prefomatted text option’.

2

I think I may have solved the problem already. I did define my spm.myTexture wrong. I am still not sure if this is the best way to deal with textures in OptiX 7 - I have a suspicion that it is not. However in the interest of potentially helping someone out in a similar situation, I am pasting my solution below.

In header file system.h:

struct Params {
texture <float4, cudaTextureType3D, cudaReadModeElementType> myTexture;
};

In myOptixTexture.cu:

include <optix.h>
include “system.h”
extern "C" __constant__ Params spm;

extern "C" __global__ void __closesthit_ch() {
// get data for T (float3)
float4 s1;
s1 = tex3D(spm.myTexture,T.x,T.y,T.z);
}

3

FWIW, in my own code I stick a cudaTextureObject_t in my launch params. You create one of these using cudaCreateTextureObject.

–
David.

4

Thanks for information @dhart.

I will check out the cudaTextureObject_t that you did in your code.

Do most people using OptiX stay away from CUDA tex3D calls?

5

Do most people using OptiX stay away from CUDA tex3D calls?

Oh, I don’t think so. Do you mean as opposed to tex2D, or something else? There is nothing at all wrong or unusual with using a CUDA tex3D call. You should feel confident in using whatever CUDA texture features you need. Don’t hesitate to call tex3D if a 3d texture is what makes sense for your application.

–
David.

6

First, in the old OptiX API there was an inconsistent support of texture samplers and OptiX bindless texture IDs.

Only the three native CUDA texture fetch calls tex1D, tex2D, and tex3D were supported on texture samplers. All others were not handled!

The way to get all other texture fetch instructions was via the OptiX texture intrinsics starting with the rtTex prefix and using bindless texture IDs.
Find all these intrinsics inside the OptiX SDK header optix_device.h.
That is the only way to access cubemaps, or have lod and gradient arguments!

OptiX bindless texture IDs also allow to check for the presence of a texture dynamically.
If the bindless texture ID is set to the predefined RT_TEXTURE_ID_NULL ( == 0) there is no texture present.
This was not possible with the other method. There the texture sampler variable always needed to be set.

In summary, it’s recommended to always use bindless texture IDs in OptiX versions before 7.0.0.

Now, OptiX 7 doesn’t know anything about textures!
There everything is handled with native CUDA calls on the host to setup CUDA texture objects and on device side with the native CUDA texture fetch functions you find in the header CUDA\<version>\include\texture_fetch_functions.h.

My OptiX applications show the comparisons between these two methods.
Old:
Creating bindless texture ID (Note that for performance in that old API, it’s recommended to build the buffer with the texture data before assigning it to the texture sampler. That’s faster when filling mipmaps into the buffer. I’m not doing that in this example.)
Declaration of the bindless texture ID in a struct
Sampling bindless texture ID

New:
Creating any CUDA texture object target (Follow the individual functions for the different texture targets in that. That later example supports mipmaps and layers.)
Declaration of the CUDA texture object in a struct
Sampling from a CUDA texture object (here 2D)

The intro_runtime example is showing CUDA texture object creation with the CUDA runtime API, all other examples there show it for the CUDA driver API. (The intro examples do not handle layers.)

1 Like
7

Thank you @droettger for the professional and in depth information. Very helpful.