在C++ class 中使用CUDA（包含texutre 2d的使用）

一直想写一下global内存访问的文章，但是最近的事情太多，哎昨天。。。，今天的一篇paper也被拒了~~哎--
不过收到一个QQ群朋友（ from this AVerMedia TECHNOLOGIES(Suzhou), Inc. ）求助邮，看得出来项目比较棘手，应该是对CUDA的使用有些不清楚；
晚上熬夜写了一个简单的C++ class 封装CUDA的demo，里面涉及到了texture的使用，希望对他的CUDA学习有帮助：
下面是三部分的代码
1. 调用class的sample.cu 文件这里没多少好解释的，就是直接生产一个cuda的class实例，然后调用各个方法

1. /********************************************************************
2. *  sample.cu
3. *  This is a example of the CUDA program.
4. *  author: zhao.kaiyong(at)gmail.com
5. *  http://www.comp.hkbu.edu.hk/~kyzhao/
6. *********************************************************************/
8. #include <stdio.h>
9. #include <stdlib.h>
10. #include <cuda_runtime.h>
11. #include <cutil.h>
12. #include "cuda_class.h"
14. /************************************************************************/
15. /* HelloCUDA                                                            */
16. /************************************************************************/
17. int main(int argc, char* argv[])
18. {
19.     cuda_class cudaA;
21.     if(!cudaA.InitCUDA()) {
22.         return 0;
23.     }
25.     float   host_data[22] = {0};
26.     float   host_result[11] ={0};
28.     for (int i = 0; i < 22; i++)
29.     {
30.         host_data[i] = i;
31.     }
33.     cudaA.InitTexture(host_data, 11, 2);
34.     cudaA.MallocMemA(11);
36.     unsigned int timer = 0;
37.     CUT_SAFE_CALL( cutCreateTimer( &timer));
38.     CUT_SAFE_CALL( cutStartTimer( timer));
40.     cudaA.DoWork();
42.     CUDA_SAFE_CALL( cudaThreadSynchronize() );
43.     CUT_SAFE_CALL( cutStopTimer( timer));
44.     printf("Processing time: %f (ms)/n", cutGetTimerValue( timer));
45.     CUT_SAFE_CALL( cutDeleteTimer( timer));
47.     cudaA.TranslateResult(host_result);
48.     cudaA.ReleaseMem();
50.     for (int i = 0; i < 11; i++)
51.     {
52.         printf("%f /n", host_result[i]);
53.     }
57.     CUT_EXIT(argc, argv);
59.     return 0;
60. }

2. 两个class 文件

1. /********************************************************************
2. *  cuda_class.h
3. *  This is a example of the CUDA program.
4. *  author: zhao.kaiyong(at)gmail.com
5. *  http://www.comp.hkbu.edu.hk/~kyzhao/
6. *********************************************************************/
7. #ifndef __CUDA_CLASS_H__
8. #define __CUDA_CLASS_H__
11. #include <cutil.h>
13. class cuda_class
14. {
15. public:
16.     cuda_class(void);
17.     ~cuda_class(void);
19.     int InitCUDA();
20.     int MallocMemA(int len);
21.     int InitTexture(float* init_data,int w, int h);
22.     int DoWork();
24.     int TranslateResult(float* out_data);
26.     int ReleaseMem();
27.     int ReleaseTex();
29. private:
30.     float   *device_result;
31.     cudaArray* device_tex;
33.     int m_ret_len;
34. };
36. #endif // __CUDA_CLASS_H__

1. /********************************************************************
2. *  cuda_class.cu
3. *  This is a example of the CUDA program.
4. *  author: zhao.kaiyong(at)gmail.com
5. *  http://www.comp.hkbu.edu.hk/~kyzhao/
6. *********************************************************************/
7. #include "cuda_class.h"
8. #include <stdio.h>
9. #include <stdlib.h>
10. #include <cuda_runtime.h>
11. #include <cutil.h>
13. texture<float, 2, cudaReadModeElementType> tex;
15. /************************************************************************/
16. /* Example                                                              */
17. /************************************************************************/
18. __global__ static void HelloCUDA(float* result, int num)
19. {
20.     int i = 0;
21.     for(i = 0; i < num; i++) {
22.         result[i] = tex2D(tex,(float) i,0) + tex2D(tex,(float)i,1);
23.     }
24. }
26. cuda_class::cuda_class(void)
27. {
28. }
30. cuda_class::~cuda_class(void)
31. {
32. }
34. int cuda_class::InitCUDA()
35. {
36.     /************************************************************************/
37.     /* Init CUDA                                                            */
38.     /************************************************************************/
39. #if __DEVICE_EMULATION__
41.     return true;
43. #else
45.     int count = 0;
46.     int i = 0;
48.     cudaGetDeviceCount(&count);
49.     if(count == 0) {
50.         fprintf(stderr, "There is no device./n");
51.         return false;
52.     }
54.     for(i = 0; i < count; i++) {
55.         cudaDeviceProp prop;
56.         if(cudaGetDeviceProperties(∝, i) == cudaSuccess) {
57.             if(prop.major >= 1) {
58.                 break;
59.             }
60.         }
61.     }
62.     if(i == count) {
63.         fprintf(stderr, "There is no device supporting CUDA./n");
64.         return false;
65.     }
66.     cudaSetDevice(i);
68.     printf("CUDA initialized./n");
69.     return true;
71. #endif
72. }
74. int cuda_class::MallocMemA(int len)
75. {
76.     m_ret_len = len;
77.     CUDA_SAFE_CALL( cudaMalloc((void**) &device_result, sizeof(float) * m_ret_len));
78.     return 1;
79. }
81. int cuda_class::DoWork()
82. {
83.     HelloCUDA<<<1, 1, 0>>>(device_result, m_ret_len);
84.     CUT_CHECK_ERROR("Kernel execution failed/n");
85.     return 1;
86. }
89. int cuda_class::TranslateResult(float * out_data)
90. {
91.     CUDA_SAFE_CALL( cudaMemcpy(out_data, device_result, sizeof(float) * m_ret_len, cudaMemcpyDeviceToHost));
92.     return 1;
93. }
95. int cuda_class::ReleaseMem()
96. {
97.     CUDA_SAFE_CALL( cudaFree(device_result));
98.     CUDA_SAFE_CALL( cudaFreeArray(device_tex));
99.     CUDA_SAFE_CALL( cudaUnbindTexture(tex));
100.     return 1;
101. }
103. int cuda_class::InitTexture(float* init_data, int w, int h)
104. {
105.     cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
107.     CUDA_SAFE_CALL( cudaMallocArray(&device_tex, &channelDesc, w, h));
108.     CUDA_SAFE_CALL( cudaMemcpyToArray(device_tex, 0, 0, init_data, sizeof(float)* w*h , cudaMemcpyHostToDevice));
109.     CUDA_SAFE_CALL( cudaBindTextureToArray(tex, device_tex, channelDesc));
110.     return 1;
111. }

这里需要做个解释，现在CUDA的cu文件是可以直接按照c++方式编译的，这个得在nvcc的编译选项里面可以设定为--host-compilation C++,在默认的情况下，以前2.0版本以前都是按照C格式编译文件，现在都是会按照C++方式编译文件，所以这里可以直接用cu文件来写Class。

在编译的时候，需要加上/MTd（debug), /MT(Release)版本，这里是告诉xcompiler选项，支持C++的runtime库，不然会有一些lib 访问冲突。

下面是下载整个工程的下载连接：不过release版本编译选项没设置，这就留给大家自己熟悉环境吧~~不要问我怎么设置……
http://download.youkuaiyun.com/user/OpenHero
cuda_cpp_class_texture_demo.rar

应该对C++和texture比较迷茫的朋友有一些帮忙.


rel="File-List" href="file:///C:%5CDOCUME%7E1%5Ckelvin%5CLOCALS%7E1%5CTemp%5Cmsohtmlclip1%5C01%5Cclip_filelist.xml"> rel="themeData" href="file:///C:%5CDOCUME%7E1%5Ckelvin%5CLOCALS%7E1%5CTemp%5Cmsohtmlclip1%5C01%5Cclip_themedata.thmx"> rel="colorSchemeMapping" href="file:///C:%5CDOCUME%7E1%5Ckelvin%5CLOCALS%7E1%5CTemp%5Cmsohtmlclip1%5C01%5Cclip_colorschememapping.xml">