Texture fetches

Integer fetch into a linear resource: no addressing, no filtering (real tex1Dfetch is the same).

Defined in code/simulator.h.

This page documents 7 symbols: 7 functions.

Functions

Signature	Description	Availability	Location	Example
if constexpr (std::is_floating_point<T>::value)	Linear filtering is modelled only for a float return type (real hardware also only filters into floats); everything else is point-sampled.	sim only	simulator.h:3167	View
inline T tex1D(cudaTextureObject_t obj, float x)		sim only	simulator.h:3155
inline T tex1Dfetch(cudaTextureObject_t obj, int x)	Integer fetch into a linear resource: no addressing, no filtering (real tex1Dfetch is the same).	sim only	simulator.h:3117
inline T tex2D(cudaTextureObject_t obj, float x, float y)		sim only	simulator.h:3159	View
inline T tex3D(cudaTextureObject_t obj, float x, float y, float z)		sim only	simulator.h:3183	View
inline void texDenorm(const TexObj& t, float& fx, float& fy, float& fz, int dims)	Denormalise the caller’s coordinates into texel space if the texture object was created with normalizedCoords.	sim only	simulator.h:3140
inline T texPoint(const TexObj& t, float fx, float fy, float fz, int dims)		sim only	simulator.h:3148

Examples

Code from the repository’s example corpus (code\*.cu); the highlighted line is the call site. These blocks are what the View links above open.

constexpr (test_llmc_matmul_backward_bias.cu:392)

    for (int k = block_d; k < x128::size; k += blockDim.z) {
        float a = 0.f;
        for (int r = warp_d; r < blockDim.z; r += bdx) {
            float v = sub_results[k][r][warp_c];
            v += __shfl_down_sync(0xffffffff, v, 1, 4);
            v += __shfl_down_sync(0xffffffff, v, 2, 4);
            a += v;
        }
        if(warp_d == 0 && global_oc < OC) {
            // coalesced, but not cacheline-sized
            if constexpr (!Atomic) {
                dbias[global_oc + k] = (OutFloat)(a + (float)dbias[global_oc + k]);
            } else {
                atomicAdd(dbias + global_oc + k, a);
            }
        }
    }
}

// Like kernel 8, but instead of accumulating to the auxiliary buffer, it writes
// multiple values that need to be summed up in a separate kernel call.
// If UseAuxBuffer is false, gridDim.y has to be one, and results are added directly

tex2D (test_web_cs_boxfilter.cu:142)

// ---- texture row pass (verbatim) --------------------------------------------
// texture fetches automatically clamp to edge of image
__global__ void d_boxfilter_x_tex(float* od, int w, int h, int r, cudaTextureObject_t tex) {
    float        scale = 1.0f / (float)((r << 1) + 1);
    unsigned int y     = blockIdx.x * blockDim.x + threadIdx.x;

    float t = 0.0f;

    for (int x = -r; x <= r; x++) {
        t += tex2D<float>(tex, (float)x, (float)y);
    }

    od[y * w] = t * scale;

    for (int x = 1; x < w; x++) {
        t += tex2D<float>(tex, (float)(x + r), (float)y);
        t -= tex2D<float>(tex, (float)(x - r - 1), (float)y);
        od[y * w + x] = t * scale;
    }
}

tex3D (test_web_cs_simpletexture3d.cu:63)

// ---- verbatim kernel from the sample ---------------------------------------
__global__ void d_render(uint* d_output, uint imageW, uint imageH, float w,
                         cudaTextureObject_t texObj) {
    uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
    uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;

    float u = x / (float)imageW;
    float v = y / (float)imageH;
    // read from 3D texture
    float voxel = tex3D<float>(texObj, u, v, w);

    if ((x < imageW) && (y < imageH)) {
        // write output color
        uint i      = __umul24(y, imageW) + x;
        d_output[i] = voxel * 255;
    }
}

// ---- host setup, from the sample's initCuda() -------------------------------
static void initCuda(const VolumeType* h_volume, cudaExtent volumeSize) {
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<VolumeType>();