CUDA 3D, 2D data processing.

In CUDA, kernels can only operate out of device memory.

Device memory can be allocated either as linear memory or CUDA arrays.

 

In this sample, we observe how to use 3D and 2D linear memory in CUDA.

Linear memory exists on the device in a 32-bit address space.

 

Typically, following methods are used for device memory allocaton.

cudaError_t cudaMalloc (void** devPtr,size_t size)

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

cudaErrot_t cudaMalloc3D (struct cudaPitchedPtr* pitchedDevPtr, struct cudaExtent extent)

 

In order to communicate host and device, following methods are used.

cudaErrot_t cudaMemcpy(void* dst, const void* src, size_t count, enum cudaMemcpyKind kind)

cudaErrot_t cudaMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t ehight, enum cudaMemcpyKind kind)

cudaError_t cudaMemcpy3D(const struct cudaMemcpy3DParms* p)

 

Following example shows how to use 3d and 2d linear memory in CUDA.

This example calculate the sum of each row, so 3D data is used input data and 2D data is result. 

1. #include <stdlib.h>
2. #include <stdio.h>
3. #include <string.h>
4. #include <math.h>
5. #include <conio.h>
6. #include <cutil_inline.h>
7. 
   
8. //-----------------------------------------------------------
9. // Device code
10. __global__ void SumRow(cudaPitchedPtr devPitchedPtr, cudaExtent extent, float* devPtr_out,int pitch_out)
11. {
12. char* devPtr   = (char*)devPitchedPtr.ptr;
13. int pitch      = devPitchedPtr.pitch;
14. int slicePitch = pitch*extent.height;
15. 
    
16. //blockDim.x  : depth
17. //threadIdx.x : height
18. char* slice = devPtr + blockIdx.x * slicePitch;
19. float* row  = (float*)(slice + threadIdx.x * pitch);
20. float* sum  = (float*)((char*)devPtr_out + blockIdx.x*pitch_out);
21. 
    
22. sum[threadIdx.x] = 0.f;
23. for(int x = 0; x<extent.width;x++){
24. sum[threadIdx.x] += row[x];
25. }
26. }
27. 
    
28. //-----------------------------------------------------------
29. // host code
30. 
    
31. int main_impl();
32. 
    
33. int main( int argc, char** argv)  {
34.  
35. main_impl();
36. }
37. 
    
38. int main_impl(){
39. 
    
40. //print_data();
41. getch();
42. size_t s[3] = {256,256,209};
43. size_t N    = s[2]*s[1]*s[0];
44. size_t size = N * sizeof(float);
45. size_t N_out    = s[2]*s[1];
46. size_t size_out = N_out * sizeof(float);
47. printf("size[%f]GB\n",size/1024.f/1204.f/1024.f);
48. 
    
49. 
    
50. // memory allocation and init. in the host memory.
51. float* h_A;
52. h_A = (float*)malloc(size);
53. float* h_B;
54. h_B = (float*)malloc(size);
55. float* h_C;
56. h_C = (float*)malloc(size_out);
57. float* h_D;
58. h_D = (float*)malloc(size_out);
59. 
    
60. for(int i=0; i<N;i++){
61. h_A[i] = i;
62. h_B[i] = N-i;
63. }
64. 
    
65. // memory allocation for device memory
66. cudaPitchedPtr devPitchedPtr;
67. cudaExtent     extent = make_cudaExtent(s[0],s[1],s[2]);
68. cudaMalloc3D(&devPitchedPtr,extent);
69. 
    
70. float* devPtr; //2D result. (y,z)
71. size_t pitch_out ;    // This variable will be set by cudaMallocPitch;
72. cudaMallocPitch( (void**)&devPtr,&pitch_out, (size_t)(s[1]*sizeof(float)), s[2]);
73. 
    
74. 
    
75. unsigned int timer = 0;
76.     cutilCheckError( cutCreateTimer( &timer));
77.     cutilCheckError( cutStartTimer( timer));
78. 
    
79. 
    
80. // copy from host memory data to device memory data
81. cudaMemcpy3DParms p = {0};
82. p.srcPtr.ptr   = h_A;
83. p.srcPtr.pitch = s[0] * sizeof(float);
84. p.srcPtr.xsize = s[0];
85. p.srcPtr.ysize = s[1];
86. p.dstPtr.ptr   = devPitchedPtr.ptr;
87. p.dstPtr.pitch = devPitchedPtr.pitch;
88. p.dstPtr.xsize = s[0];
89. p.dstPtr.ysize = s[1];
90. p.extent.width  = s[0]*sizeof(float);
91. p.extent.height = s[1];
92. p.extent.depth  = s[2];
93. p.kind = cudaMemcpyHostToDevice;
94. cudaMemcpy3D(&p);
95. 
    
96. int threadsPerBlock = s[1];
97. int blocksPerGrid   = s[2];
98. SumRow<<<blocksPerGrid,threadsPerBlock>>>(devPitchedPtr,extent,devPtr,pitch_out);
99. cudaGetErrorString(cudaGetLastError());
100. cudaMemcpy2D(h_C,s[1]*sizeof(float),devPtr,pitch_out,s[1]*sizeof(float),s[2],cudaMemcpyDeviceToHost);
101. // free device memory.
102. cudaFree(devPitchedPtr.ptr);
103. 
     
104. // finish Timer
105. cutilCheckError( cutStopTimer( timer));
106. printf( "Processing time1: %f (ms)\n", cutGetTimerValue( timer));
107. cutilCheckError( cutDeleteTimer( timer));
108. 
     
109. 
     
110. // for reference.
111. timer = 0;
112. cutilCheckError( cutCreateTimer( &timer));
113. cutilCheckError( cutStartTimer( timer));
114. int i=0;
115. int j=0;
116. for(int z=0; z< s[2];z++){
117. for(int y=0; y< s[1];y++){
118. h_D[j] =0.f;
119. for(int x=0; x< s[0];x++){
120. h_D[j] += h_A[i];
121. i++;
122. }
123. j++;
124. }
125. }
126. cutilCheckError( cutStopTimer( timer));
127. printf( "Processing time2: %f (ms)\n", cutGetTimerValue( timer));
128. cutilCheckError( cutDeleteTimer( timer));
129. 
     
130. // comparison
131. for(i=0; i< N_out;i++){
132. 
     
133. // printf("test N[%d][%6.5f/%6.5f]\n",i,h_C[i],h_D[i]);
134. if( 0.0f < (h_C[i] - h_D[i]) ){
135. printf("error N[%d][%6.5f/%6.5f]\n",i,h_C[i],h_D[i]);
136. }
137. }
138. return 0;
139. };

 

 

 

이 글은 스프링노트에서 작성되었습니다.