Texture object state

Defined in code/simulator.h.

This page documents 15 symbols: 5 enums, 10 functions.

Enums

Signature	Description	Availability	Location	Example
enum cudaResourceType		sim only	simulator.h:2988
enum cudaSurfaceBoundaryMode		sim only	simulator.h:3023
enum cudaTextureAddressMode		sim only	simulator.h:2992
enum cudaTextureFilterMode		sim only	simulator.h:2996
enum cudaTextureReadMode		sim only	simulator.h:2997

Functions

Signature	Description	Availability	Location	Example
inline bool addrApply(cudaTextureAddressMode m, long long& i, long long n, bool normalized = true)	Apply an address mode to an integer coordinate against extent n.	sim only	simulator.h:3037
if constexpr (NormPack<T>::ok)		sim only	simulator.h:3062	View
inline cudaError_t cudaCreateSurfaceObject(cudaSurfaceObject_t* obj, const cudaResourceDesc* res)		sim only	simulator.h:3104
inline cudaError_t cudaCreateTextureObject(cudaTextureObject_t* obj, const cudaResourceDesc* res, const cudaTextureDesc* tex, const void* )		sim only	simulator.h:3094	View
inline cudaError_t cudaDestroySurfaceObject(cudaSurfaceObject_t )		sim only	simulator.h:3112
inline cudaError_t cudaDestroyTextureObject(cudaTextureObject_t )		sim only	simulator.h:3103	View
inline TexObj* texAt(cudaTextureObject_t o)		sim only	simulator.h:3029
inline T texelAt(const TexObj& t, long long x, long long y, long long z)	Raw texel read out of whatever storage the resource uses.	sim only	simulator.h:3055
inline void texExtent(const TexObj& t, long long& w, long long& h, long long& d)		sim only	simulator.h:3079
inline std::vector<TexObj>& texObjs()		sim only	simulator.h:3028

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

cudaCreateTextureObject (test_web_cs_boxfilter.cu:227)

    cudaTextureDesc texDescr;
    std::memset(&texDescr, 0, sizeof(cudaTextureDesc));
    texDescr.normalizedCoords = 0;                      // [SIM] integer texel coords
    texDescr.filterMode       = cudaFilterModePoint;    // [SIM] see header note
    texDescr.addressMode[0]   = cudaAddressModeClamp;
    texDescr.addressMode[1]   = cudaAddressModeClamp;
    texDescr.readMode         = cudaReadModeElementType;

    cudaTextureObject_t tex = 0;
    check(cudaCreateTextureObject(&tex, &texRes, &texDescr, NULL), "cudaCreateTextureObject");

    float *d_src = nullptr, *d_temp = nullptr, *d_temp2 = nullptr, *d_dest = nullptr;
    check(cudaMalloc(reinterpret_cast<void**>(&d_src),   nByte), "cudaMalloc src");
    check(cudaMalloc(reinterpret_cast<void**>(&d_temp),  nByte), "cudaMalloc temp");
    check(cudaMalloc(reinterpret_cast<void**>(&d_temp2), nByte), "cudaMalloc temp2");
    check(cudaMalloc(reinterpret_cast<void**>(&d_dest),  nByte), "cudaMalloc dest");
    check(cudaMemcpy(d_src, hImg.data(), nByte, cudaMemcpyHostToDevice), "cudaMemcpy H2D");

    // Row pass through the texture, then the column pass on global memory.
    LAUNCH(d_boxfilter_x_tex, dim3(height / nthreads), dim3(nthreads),
           d_temp, width, height, radius, tex);

cudaDestroyTextureObject (test_web_cs_boxfilter.cu:267)

    boxFilterRef(hImg, ref, width, height, radius);
    double maxRel = 0.0;
    for (size_t i = 0; i < nPix; ++i) {
        const double den = std::fabs(ref[i]) > 1e-6 ? std::fabs(ref[i]) : 1.0;
        maxRel = std::max(maxRel, std::fabs((double)hOut[i] - ref[i]) / den);
    }
    std::printf("box filter (r=%d, %dx%d) vs direct-sum reference: max rel err = %g\n",
                radius, width, height, maxRel);
    if (!(maxRel < 1e-5)) g_ok = false;

    check(cudaDestroyTextureObject(tex), "cudaDestroyTextureObject");
    check(cudaFreeArray(d_array), "cudaFreeArray");
    check(cudaFree(d_src),   "cudaFree src");
    check(cudaFree(d_temp),  "cudaFree temp");
    check(cudaFree(d_temp2), "cudaFree temp2");
    check(cudaFree(d_dest),  "cudaFree dest");

    std::printf("boxFilter: %s\n", g_ok ? "PASSED" : "FAILED");
    return g_ok ? 0 : 1;
}