3-D arrays#

Defined in code/simulator.h.

This page documents 9 symbols: 9 functions.

Functions#

Signature

Description

Availability

Location

Example

inline cudaError_t cudaGetChannelDesc(cudaChannelFormatDesc* d, cudaArray_const_t a)

sim only

simulator.h:2957

inline cudaError_t cudaMalloc3DArray(cudaArray_t* a, const cudaChannelFormatDesc* desc, cudaExtent ext, unsigned = 0)

sim only

simulator.h:2898

View

inline cudaError_t cudaMemcpy2DFromArray(void* dst, size_t dpitch, cudaArray_const_t src, size_t wOffset, size_t hOffset, size_t width, size_t height, cudaMemcpyKind )

sim only

simulator.h:2975

inline cudaError_t cudaMemcpy2DToArray(cudaArray_t dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, cudaMemcpyKind )

sim only

simulator.h:2960

View

inline cudaError_t cudaMemcpy3D(const cudaMemcpy3DParms* p)

Copies a slab between a linear pitched buffer and a 3-D cudaArray.

sim only

simulator.h:2923

View

inline cudaError_t cudaMemcpyToArray(cudaArray_t dst, size_t wOffset, size_t hOffset, const void* src, size_t count, cudaMemcpyKind kind)

sim only

simulator.h:2971

View

inline cudaExtent make_cudaExtent(size_t w, size_t h, size_t d)

sim only

simulator.h:2891

View

inline cudaPitchedPtr make_cudaPitchedPtr(void* p, size_t pitch, size_t xsz, size_t ysz)

sim only

simulator.h:2895

View

inline cudaPos make_cudaPos(size_t x, size_t y, size_t z)

sim only

simulator.h:2893

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.

cudaMalloc3DArray (test_web_cs_simpletexture3d.cu:75)
65    if ((x < imageW) && (y < imageH)) {
66        // write output color
67        uint i      = __umul24(y, imageW) + x;
68        d_output[i] = voxel * 255;
69    }
70}
71
72// ---- host setup, from the sample's initCuda() -------------------------------
73static void initCuda(const VolumeType* h_volume, cudaExtent volumeSize) {
74    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<VolumeType>();
75    check(cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumeSize), "cudaMalloc3DArray");
76
77    cudaMemcpy3DParms copyParams = {};
78    copyParams.srcPtr   = make_cudaPitchedPtr((void*)h_volume,
79                                              volumeSize.width * sizeof(VolumeType),
80                                              volumeSize.width, volumeSize.height);
81    copyParams.dstArray = d_volumeArray;
82    copyParams.extent   = make_cudaExtent(volumeSize.width, volumeSize.height, volumeSize.depth);
83    copyParams.kind     = cudaMemcpyHostToDevice;
84    check(cudaMemcpy3D(&copyParams), "cudaMemcpy3D");
85
86    cudaResourceDesc texRes;
cudaMemcpy2DToArray (test_web_cs_boxfilter.cu:209)
199    std::mt19937 rng(12345u);
200    std::uniform_int_distribution<int> d16(0, 15);
201    std::vector<float> hImg(nPix);
202    for (size_t i = 0; i < nPix; ++i) hImg[i] = (float)d16(rng);
203
204    // Source image in a cudaArray, plus a plain linear copy for the non-texture
205    // row pass.
206    cudaChannelFormatDesc fdesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
207    cudaArray_t d_array = nullptr;
208    check(cudaMallocArray(&d_array, &fdesc, width, height), "cudaMallocArray");
209    check(cudaMemcpy2DToArray(d_array, 0, 0, hImg.data(), width * sizeof(float),
210                              width * sizeof(float), height, cudaMemcpyHostToDevice),
211          "cudaMemcpy2DToArray");
212
213    cudaResourceDesc texRes;
214    std::memset(&texRes, 0, sizeof(cudaResourceDesc));
215    texRes.resType         = cudaResourceTypeArray;
216    texRes.res.array.array = d_array;
217
218    cudaTextureDesc texDescr;
219    std::memset(&texDescr, 0, sizeof(cudaTextureDesc));
220    texDescr.normalizedCoords = 0;                      // [SIM] integer texel coords
cudaMemcpy3D (test_web_cs_simpletexture3d.cu:84)
74    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<VolumeType>();
75    check(cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumeSize), "cudaMalloc3DArray");
76
77    cudaMemcpy3DParms copyParams = {};
78    copyParams.srcPtr   = make_cudaPitchedPtr((void*)h_volume,
79                                              volumeSize.width * sizeof(VolumeType),
80                                              volumeSize.width, volumeSize.height);
81    copyParams.dstArray = d_volumeArray;
82    copyParams.extent   = make_cudaExtent(volumeSize.width, volumeSize.height, volumeSize.depth);
83    copyParams.kind     = cudaMemcpyHostToDevice;
84    check(cudaMemcpy3D(&copyParams), "cudaMemcpy3D");
85
86    cudaResourceDesc texRes;
87    std::memset(&texRes, 0, sizeof(cudaResourceDesc));
88    texRes.resType         = cudaResourceTypeArray;
89    texRes.res.array.array = d_volumeArray;
90
91    cudaTextureDesc texDescr;
92    std::memset(&texDescr, 0, sizeof(cudaTextureDesc));
93    texDescr.normalizedCoords = true;              // access with normalized coords
94    texDescr.filterMode       = cudaFilterModeLinear;   // linear interpolation
95    texDescr.addressMode[0]   = cudaAddressModeWrap;    // wrap texture coordinates
cudaMemcpyToArray (test_web_cs_convtexture.cu:229)
219    texDescr.normalizedCoords = 0;
220    texDescr.filterMode       = cudaFilterModeLinear;
221    texDescr.addressMode[0]   = cudaAddressModeClamp;   // [SIM] sample says Wrap
222    texDescr.addressMode[1]   = cudaAddressModeClamp;   //       (see header note)
223    texDescr.readMode         = cudaReadModeElementType;
224
225    cudaTextureObject_t texSrc = 0;
226    check(cudaCreateTextureObject(&texSrc, &texRes, &texDescr, NULL), "cudaCreateTextureObject");
227
228    setConvolutionKernel(h_Kernel.data());
229    check(cudaMemcpyToArray(a_Src, 0, 0, h_Input.data(), nByte, cudaMemcpyHostToDevice),
230          "cudaMemcpyToArray H2D");
231
232    convolutionRowsGPU(d_Output, a_Src, imageW, imageH, texSrc);
233    check(cudaDeviceSynchronize(), "sync after rows");
234
235    // Kernels cannot write to a texture, so the row result goes back to the array.
236    check(cudaMemcpyToArray(a_Src, 0, 0, d_Output, nByte, cudaMemcpyDeviceToDevice),
237          "cudaMemcpyToArray D2D");
238
239    convolutionColumnsGPU(d_Output, a_Src, imageW, imageH, texSrc);
240    check(cudaDeviceSynchronize(), "sync after columns");
make_cudaExtent (test_web_cs_simpletexture3d.cu:82)
72// ---- host setup, from the sample's initCuda() -------------------------------
73static void initCuda(const VolumeType* h_volume, cudaExtent volumeSize) {
74    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<VolumeType>();
75    check(cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumeSize), "cudaMalloc3DArray");
76
77    cudaMemcpy3DParms copyParams = {};
78    copyParams.srcPtr   = make_cudaPitchedPtr((void*)h_volume,
79                                              volumeSize.width * sizeof(VolumeType),
80                                              volumeSize.width, volumeSize.height);
81    copyParams.dstArray = d_volumeArray;
82    copyParams.extent   = make_cudaExtent(volumeSize.width, volumeSize.height, volumeSize.depth);
83    copyParams.kind     = cudaMemcpyHostToDevice;
84    check(cudaMemcpy3D(&copyParams), "cudaMemcpy3D");
85
86    cudaResourceDesc texRes;
87    std::memset(&texRes, 0, sizeof(cudaResourceDesc));
88    texRes.resType         = cudaResourceTypeArray;
89    texRes.res.array.array = d_volumeArray;
90
91    cudaTextureDesc texDescr;
92    std::memset(&texDescr, 0, sizeof(cudaTextureDesc));
93    texDescr.normalizedCoords = true;              // access with normalized coords
make_cudaPitchedPtr (test_web_cs_simpletexture3d.cu:78)
68        d_output[i] = voxel * 255;
69    }
70}
71
72// ---- host setup, from the sample's initCuda() -------------------------------
73static void initCuda(const VolumeType* h_volume, cudaExtent volumeSize) {
74    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<VolumeType>();
75    check(cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumeSize), "cudaMalloc3DArray");
76
77    cudaMemcpy3DParms copyParams = {};
78    copyParams.srcPtr   = make_cudaPitchedPtr((void*)h_volume,
79                                              volumeSize.width * sizeof(VolumeType),
80                                              volumeSize.width, volumeSize.height);
81    copyParams.dstArray = d_volumeArray;
82    copyParams.extent   = make_cudaExtent(volumeSize.width, volumeSize.height, volumeSize.depth);
83    copyParams.kind     = cudaMemcpyHostToDevice;
84    check(cudaMemcpy3D(&copyParams), "cudaMemcpy3D");
85
86    cudaResourceDesc texRes;
87    std::memset(&texRes, 0, sizeof(cudaResourceDesc));
88    texRes.resType         = cudaResourceTypeArray;
89    texRes.res.array.array = d_volumeArray;
On this page