Pitched allocation / 2-D copy#

Defined in code/simulator.h.

This page documents 2 symbols: 2 functions.

Functions#

Signature

Description

Availability

Location

Example

inline cudaError_t cudaMallocPitch(void** devPtr, size_t* pitch, size_t width, size_t height)

sim only

simulator.h:2367

View

inline cudaError_t cudaMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, cudaMemcpyKind )

sim only

simulator.h:2375

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.

cudaMallocPitch (test_web_cs_dct8x8.cu:467)
457    double rtErr = 0.0;
458    for (size_t i = 0; i < nPix; ++i) rtErr = std::max(rtErr, std::fabs(goldRound[i] - ImgD[i]));
459    std::printf("gold DCT->IDCT round trip: max |diff| = %g\n", rtErr);
460
461    // ---- device buffers -----------------------------------------------------
462    cudaChannelFormatDesc floattex = cudaCreateChannelDesc<float>();
463    cudaArray* Src = nullptr;
464    float*     Dst = nullptr;
465    size_t     DstStride = 0;
466    check(cudaMallocArray(&Src, &floattex, w, h), "cudaMallocArray");
467    check(cudaMallocPitch(reinterpret_cast<void**>(&Dst), &DstStride, w * sizeof(float), h),
468          "cudaMallocPitch");
469    const size_t DstStrideBytes = DstStride;
470    DstStride /= sizeof(float);
471
472    check(cudaMemcpy2DToArray(Src, 0, 0, ImgF.data(), w * sizeof(float), w * sizeof(float), h,
473                              cudaMemcpyHostToDevice), "cudaMemcpy2DToArray H2A");
474
475    cudaResourceDesc texRes;
476    std::memset(&texRes, 0, sizeof(cudaResourceDesc));
477    texRes.resType         = cudaResourceTypeArray;
478    texRes.res.array.array = Src;
cudaMemcpy2D (test_web_cs_dct8x8.cu:499)
489    check(cudaCreateTextureObject(&TexSrc, &texRes, &texDescr, NULL), "cudaCreateTextureObject");
490
491    dim3 threads(BLOCK_SIZE, BLOCK_SIZE);
492    dim3 grid(w / BLOCK_SIZE, h / BLOCK_SIZE);
493
494    // ---- kernel 1: DCT -> quantize -> (back to the array) -> IDCT -----------
495    LAUNCH(CUDAkernel1DCT, grid, threads, Dst, (int)DstStride, 0, 0, TexSrc);
496    check(cudaDeviceSynchronize(), "sync k1 DCT");
497
498    std::vector<float> coef1(nPix, 0.0f);
499    check(cudaMemcpy2D(coef1.data(), w * sizeof(float), Dst, DstStrideBytes, w * sizeof(float), h,
500                       cudaMemcpyDeviceToHost), "D2H coef1");
501
502    LAUNCH(CUDAkernelQuantizationFloat, grid, threads, Dst, (int)DstStride);
503    check(cudaDeviceSynchronize(), "sync k1 quant");
504
505    check(cudaMemcpy2DToArray(Src, 0, 0, Dst, DstStrideBytes, w * sizeof(float), h,
506                              cudaMemcpyDeviceToDevice), "cudaMemcpy2DToArray D2A");
507
508    LAUNCH(CUDAkernel1IDCT, grid, threads, Dst, (int)DstStride, 0, 0, TexSrc);
509    check(cudaDeviceSynchronize(), "sync k1 IDCT");
On this page