Textures and surfaces

Textures and surfaces (sim only) A CPU model of the bindless texture API: cudaArray-backed storage, a channel descriptor, and texture objects carrying addressing + filtering state. MODELLED: 1-D/2-D/3-D fetches by integer or unnormalised/normalised float coordinates, all four address modes (wrap/clamp/mirror/border) and both filter modes; cudaReadModeNormalizedFloat for cudaArray-backed 8- and 16-bit integer texels; surfaces via surf1D/surf2Dread/write. NOT modelled: mipmaps, layered and cubemap textures, sRGB conversion and anisotropy — texture memory here is just an array with addressing rules, not a cache.

Defined in code/simulator.h.

This page documents 10 symbols: 1 enums, 9 functions.

Enums

Signature	Description	Availability	Location	Example
enum cudaChannelFormatKind		sim only	simulator.h:2788

Functions

Signature	Description	Availability	Location	Example
inline cudaChannelFormatDesc cudaCreateChannelDesc(int x, int y, int z, int w, cudaChannelFormatKind f)		sim only	simulator.h:2794	View
inline cudaChannelFormatDesc cudaCreateChannelDesc()		sim only	simulator.h:2814	View
inline cudaError_t cudaFreeArray(cudaArray_t a)		sim only	simulator.h:2887	View
inline cudaError_t cudaMallocArray(cudaArray_t* a, const cudaChannelFormatDesc* desc, size_t width, size_t height = 0, unsigned = 0)		sim only	simulator.h:2875	View
inline void decodeNormalized(const cudaChannelFormatDesc& d, const unsigned char* p, float out[4])	cudaReadModeNormalizedFloat: an 8- or 16-bit integer texel is returned as a float, unsigned scaled to [0,1] and signed to [-1,1] (the real hardware conversion).	sim only	simulator.h:2836
static cudaChannelFormatDesc get()		sim only	simulator.h:2800
static float pack(const float c[4])		sim only	simulator.h:2868	View
static float2 pack(const float c[4])		sim only	simulator.h:2870	View
static float4 pack(const float c[4])		sim only	simulator.h:2872	View

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.

cudaCreateChannelDesc (test_web_cs_boxfilter.cu:206)

    // [SIM] small integers, so the sliding-window sums stay exact and the two
    // row passes (texture vs global) can be required to match bit for bit.
    std::mt19937 rng(12345u);
    std::uniform_int_distribution<int> d16(0, 15);
    std::vector<float> hImg(nPix);
    for (size_t i = 0; i < nPix; ++i) hImg[i] = (float)d16(rng);

    // Source image in a cudaArray, plus a plain linear copy for the non-texture
    // row pass.
    cudaChannelFormatDesc fdesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
    cudaArray_t d_array = nullptr;
    check(cudaMallocArray(&d_array, &fdesc, width, height), "cudaMallocArray");
    check(cudaMemcpy2DToArray(d_array, 0, 0, hImg.data(), width * sizeof(float),
                              width * sizeof(float), height, cudaMemcpyHostToDevice),
          "cudaMemcpy2DToArray");

    cudaResourceDesc texRes;
    std::memset(&texRes, 0, sizeof(cudaResourceDesc));
    texRes.resType         = cudaResourceTypeArray;
    texRes.res.array.array = d_array;

cudaCreateChannelDesc (test_web_cs_boxfilter.cu:206)

    // [SIM] small integers, so the sliding-window sums stay exact and the two
    // row passes (texture vs global) can be required to match bit for bit.
    std::mt19937 rng(12345u);
    std::uniform_int_distribution<int> d16(0, 15);
    std::vector<float> hImg(nPix);
    for (size_t i = 0; i < nPix; ++i) hImg[i] = (float)d16(rng);

    // Source image in a cudaArray, plus a plain linear copy for the non-texture
    // row pass.
    cudaChannelFormatDesc fdesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
    cudaArray_t d_array = nullptr;
    check(cudaMallocArray(&d_array, &fdesc, width, height), "cudaMallocArray");
    check(cudaMemcpy2DToArray(d_array, 0, 0, hImg.data(), width * sizeof(float),
                              width * sizeof(float), height, cudaMemcpyHostToDevice),
          "cudaMemcpy2DToArray");

    cudaResourceDesc texRes;
    std::memset(&texRes, 0, sizeof(cudaResourceDesc));
    texRes.resType         = cudaResourceTypeArray;
    texRes.res.array.array = d_array;

cudaFreeArray (test_web_cs_boxfilter.cu:268)

    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;
}

cudaMallocArray (test_web_cs_boxfilter.cu:208)

    // row passes (texture vs global) can be required to match bit for bit.
    std::mt19937 rng(12345u);
    std::uniform_int_distribution<int> d16(0, 15);
    std::vector<float> hImg(nPix);
    for (size_t i = 0; i < nPix; ++i) hImg[i] = (float)d16(rng);

    // Source image in a cudaArray, plus a plain linear copy for the non-texture
    // row pass.
    cudaChannelFormatDesc fdesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
    cudaArray_t d_array = nullptr;
    check(cudaMallocArray(&d_array, &fdesc, width, height), "cudaMallocArray");
    check(cudaMemcpy2DToArray(d_array, 0, 0, hImg.data(), width * sizeof(float),
                              width * sizeof(float), height, cudaMemcpyHostToDevice),
          "cudaMemcpy2DToArray");

    cudaResourceDesc texRes;
    std::memset(&texRes, 0, sizeof(cudaResourceDesc));
    texRes.resType         = cudaResourceTypeArray;
    texRes.res.array.array = d_array;

    cudaTextureDesc texDescr;
    std::memset(&texDescr, 0, sizeof(cudaTextureDesc));

pack (test_web_pk_packsplit.cu:189)

    cudaDeviceSynchronize();

    cudaMemcpy(h_output_gpu.data(), d_output, sizeof(float) * length, cudaMemcpyDeviceToHost);
    // CPU reference produces a compacted prefix; unused tail is undefined, so
    // only the first (count of positives) entries are compared.
    std::vector<float> h_output_pack(length, 0.f);
    pack_host(h_output_pack.data(), h_input.data(), length);
    int n_pos = 0;
    for (int i = 0; i < length; i++) if (h_input[i] > FLT_ZERO) n_pos++;
    bool pack_ok = validation(h_output_pack.data(), h_output_gpu.data(), n_pos);
    std::printf("  pack  (%d elems, %d positive): %s\n", length, n_pos, pack_ok ? "match" : "MISMATCH");
    ok = ok && pack_ok;

    // ---- Split --------------------------------------------------------------
    cudaMemcpy(d_input, d_output, sizeof(float) * length, cudaMemcpyDeviceToDevice);
    cudaMemset(d_output, 0, sizeof(float) * length);
    LAUNCH(split_kernel, GRID_DIM, N_ELEM, d_output, d_input, d_predicates, d_scanned, length);  // [SIM]
    cudaDeviceSynchronize();

    cudaMemcpy(h_output_gpu.data(), d_output, sizeof(float) * length, cudaMemcpyDeviceToHost);
    split_host(h_output_host.data(), h_input.data(), length);
    bool split_ok = validation(h_output_host.data(), h_output_gpu.data(), length);

pack (test_web_pk_packsplit.cu:189)

    cudaDeviceSynchronize();

    cudaMemcpy(h_output_gpu.data(), d_output, sizeof(float) * length, cudaMemcpyDeviceToHost);
    // CPU reference produces a compacted prefix; unused tail is undefined, so
    // only the first (count of positives) entries are compared.
    std::vector<float> h_output_pack(length, 0.f);
    pack_host(h_output_pack.data(), h_input.data(), length);
    int n_pos = 0;
    for (int i = 0; i < length; i++) if (h_input[i] > FLT_ZERO) n_pos++;
    bool pack_ok = validation(h_output_pack.data(), h_output_gpu.data(), n_pos);
    std::printf("  pack  (%d elems, %d positive): %s\n", length, n_pos, pack_ok ? "match" : "MISMATCH");
    ok = ok && pack_ok;

    // ---- Split --------------------------------------------------------------
    cudaMemcpy(d_input, d_output, sizeof(float) * length, cudaMemcpyDeviceToDevice);
    cudaMemset(d_output, 0, sizeof(float) * length);
    LAUNCH(split_kernel, GRID_DIM, N_ELEM, d_output, d_input, d_predicates, d_scanned, length);  // [SIM]
    cudaDeviceSynchronize();

    cudaMemcpy(h_output_gpu.data(), d_output, sizeof(float) * length, cudaMemcpyDeviceToHost);
    split_host(h_output_host.data(), h_input.data(), length);
    bool split_ok = validation(h_output_host.data(), h_output_gpu.data(), length);

pack (test_web_pk_packsplit.cu:189)

    cudaDeviceSynchronize();

    cudaMemcpy(h_output_gpu.data(), d_output, sizeof(float) * length, cudaMemcpyDeviceToHost);
    // CPU reference produces a compacted prefix; unused tail is undefined, so
    // only the first (count of positives) entries are compared.
    std::vector<float> h_output_pack(length, 0.f);
    pack_host(h_output_pack.data(), h_input.data(), length);
    int n_pos = 0;
    for (int i = 0; i < length; i++) if (h_input[i] > FLT_ZERO) n_pos++;
    bool pack_ok = validation(h_output_pack.data(), h_output_gpu.data(), n_pos);
    std::printf("  pack  (%d elems, %d positive): %s\n", length, n_pos, pack_ok ? "match" : "MISMATCH");
    ok = ok && pack_ok;

    // ---- Split --------------------------------------------------------------
    cudaMemcpy(d_input, d_output, sizeof(float) * length, cudaMemcpyDeviceToDevice);
    cudaMemset(d_output, 0, sizeof(float) * length);
    LAUNCH(split_kernel, GRID_DIM, N_ELEM, d_output, d_input, d_predicates, d_scanned, length);  // [SIM]
    cudaDeviceSynchronize();

    cudaMemcpy(h_output_gpu.data(), d_output, sizeof(float) * length, cudaMemcpyDeviceToHost);
    split_host(h_output_host.data(), h_input.data(), length);
    bool split_ok = validation(h_output_host.data(), h_output_gpu.data(), length);