大家好：
我是刚刚研究GPU不久的，遇到一个问题，希望牛人可以帮忙：
目前我在单机单GPU上使用多线程调度GPU，就是CPU上两个线程，一个线程调用Kernel函数，一个线程执行CPU计算，这样的结果理论上不应该是比单线程调用（先调用Kernel函数，执行完后再执行CPU计算函数）结果速度快吗？因为GPU计算和CPU计算可以并发执行，可是结果是几乎没变化，有时甚至比单线程还慢。PS:CPU是多核的
貌似时间短的程序效果不是很明显，换了时间长的程序后有些加速效果
LZ你好，一般CUDA的内核在其发射时就已经返回返回值给CPU端，故一般在你未使用显示或隐式同步的情况下，内核执行和内核后面CPU代码执行本来就是同时的。
如：
kernel<<<…>>>(…);
hostcode1(…);
cudaDeviceSynchronize();
hostcode2(…);
假设kernel的执行时间大于hostcode1+hostcode2的执行时间，按照上面的写法，kernel是和hostcode1一同执行的，由于执行了cudaDeviceSyncthronize同步，hostcode2是在kernel执行完成之后才执行的。
如果你的程序只是像上例中前两行的写法，那样你是否应用多线程，并没有什么本质区别。
谢谢回复。
上面的cudaDeviceSynchronize();是我单线程的写法，多线程情况下控制Kernel的线程使用同步函数，并不影响控制Host函数的线程执行情况。
楼主您好，
一个代码快不快，是一个客观的问题，也是您的个人看法/感觉的问题，
如果不考虑后者，则在同样的硬件下，更多的取决与您的程序的编写水平。
论坛无法提供一般的通用的理论性的参考能让您直接用到您的程序中，立刻让您提高速度；
但您可以通过提供您的具体程序代码，来让论坛帮助您看下是否有可以改进的地方。
请理解并谅解此点。
谢谢版主的提醒，我把我写的代码粘上来，希望大家给我点建议，谢谢哈！
35 void *thread1(void *){
36 float time, start;
37 start = clock();
38 printf(“thread1: begin\n”);
39 int i,n=10000;
40 float *a,*b,*c;
41 a = (float )malloc(nsizeof(float));
42 b = (float )malloc(nsizeof(float));
43 c = (float )malloc(nsizeof(float));
44 for(i = 0; i <n;i++){
45 a[i] = 1.0f;
46 b[i] = 1.0f;
47 }
48 // for(i = 0; i < 50000; i++){
49 vectorAdd(a,b,c,n);
50 // }
51 // cudaThreadSynchronize();
52 printf(“thread1:c[%d]=%f\n”,0,c[0]);
53 free(a);
54 free(b);
55 free(c);
56 time = clock()-start;
57 printf(“thread1: end, and execute time is :%f\n”, time/CLOCKS_PER_SEC);
58 pthread_exit(NULL);
59 }
61 void *thread2(void *){
62 float time, start;
63
64 start = clock();
65
66 printf(“thread2:I’m thread2n”);
67
68 int i, j, k = 1;
69
70 for(i = 0; i < 500000; i++)
71 {
72 for(j = 0; j < 100000; j++)
73 {
74 // printf(“thread2:k = %d\n”, k);
75 k++;
76 }
77 }
78 time= clock() - start;
79 printf(“thread2: task was finished!ncostTime2 : %f\n”, time / CLOCKS_PER_SEC);
80 pthread_exit(NULL);
81 }
多线程测试里面使用 pthread_join(thread[0], NULL);pthread_join(thread[1], NULL);
单线程测试里面就直接调用上面两个函数，中间执行一次cudaThreadSynchronize();
单线程测试结果是106ms，多线程测试结果是94.75ms，这是在将循环次数调到很大的时候，如果调小一点基本上就没有什么效果了
楼主您好，您的代码有几个严重问题。
（1）因为您的kernel的指针使用，直接使用了host上的malloc出来的值，导致它根本不会执行。而您“测试”出来的时间，只是runtime初始化+host code的执行时间。而完全不涉及kernel.
（2）您的host上的对时间的测试是完全错误的，您使用了clock()函数，这个函数在NIX上将不是生活中的时钟。我具体说下：
（2-1）clock()在NIX上将只会测试进程的cpu时间片的和。而不是实际时间。
（2-1）clock()在*NIX上，如果使用多个核心，时间将是多个的累加（而不是您想象中的算一份，例如2个核心都100%的运行了实际生活中的时间10s, 将被计算为2份共20s）。
所以第二点您可以看出，只有一个进程100%的占用CPU的一个核心，用clock计时才是靠谱的。一般CPU使用率，或者核心数发生变化，这个得到的时间将对您的目的来说毫无意义。
因为您实际上从来没有执行kernel，
在外加您实际上的测时完全错误，
所以您的结论毫无意义。
请修正上文指出的2个问题。谢谢。
（您只有先在逻辑上正确，才能考虑优化之类的问题，您觉得呢？）
感谢周末来访。
谢谢版主的认真解读，我担心篇幅太长会让大家失去兴趣，所以就简要的把代码贴出来了，不好意思，可能是我没表达清楚，加上没有粘贴核函数的代码，所以让您误以为thread1就是核函数，下面我把核函数粘贴一下
pthread_t thread[2];
8 global void vectorAddKernel(float *a, float *b, float *c, int n){
9 int x = threadIdx.x + blockIdx.x * blockDim.x;
10 int y = threadIdx.y + blockIdx.y * blockDim.y;
11 int index = x + y * gridDim.x * blockDim.x;
12 if(index < n){
13 for(int i = 0; i< 50000000; i++){
14 for(int j = 0; j < 500; j++){
15 c[index] = a[index] + b[index];
16 }
17 }
18
19 }
20 }
21 void vectorAdd(float *a, float b, floatc, int n){
22 float *d_a, *d_b, *d_c;
23 cudaMalloc((void **)&d_a, n * sizeof(float));
24 cudaMemcpy(d_a, a, n * sizeof(float), cudaMemcpyHostToDevice);
25 cudaMalloc((void **)&d_b, n * sizeof(float));
26 cudaMemcpy(d_b, b, n * sizeof(float), cudaMemcpyHostToDevice);
27 cudaMalloc((void **)&d_c, n * sizeof(float));
28 cudaMemcpy(d_c, b, n * sizeof(float), cudaMemcpyHostToDevice);
29 vectorAddKernel<<<10, 1024>>>(d_a, d_b, d_c, n);
30 cudaMemcpy(c, d_c, n * sizeof(float), cudaMemcpyDeviceToHost);
31 cudaFree(d_a);
32 cudaFree(d_b);
33 cudaFree(d_c);
34 }
，其实上面的测试结果统计时间是在main函数中统计的整个程序运行完的时间：
104 float time, start;
105 start = clock();
106 thread_create();
107 printf(“main thread!!!\n”);
108 thread_wait();
109 time = clock() - start;
110 printf(“main thread cost time is :%f\n”, time/CLOCKS_PER_SEC);
因为我只想知道多线程情况下，程序总体运行时间是否会减少，所以就只考虑这个main函数中的总体时间是多少。
嗯嗯。 上面的那个的确不是你的kernel, 这个是我理解错了。
现在就您的最新代码给出回复：
因为您的2个线程之间分别是GPU的kernel处理和另外的一个毫无关系的CPU任务(printf一堆东西)，
那么显然执行时间会减少的。
我用个图表示下，假设：
thread1()代码原来需要5分钟（wall time, 实际生活中的时间，不是实际占用的CPU时间)
thread2()代码原来需要2分钟 (同上)
那么原来您在6月16日的下午18:30分开始串行执行:
thread1();
thread2();
那么显然您将在今天下午的18:37分才能看到程序结束。
但是如果您此时同时开启了2个host thread, 分别执行thread1()和thread2()这2个函数，
那么您显然在下午的18:35分就看到程序结束了。
所以结论显然，您的代码的确可以看到提前结束。
而这个和您有几个核心无关，1个核心也能观察到这个效果。甚至和您的cudaMemcpy()在隐式同步时候的是否是100% CPU占用的spin等待还是0% blocking等待也无关。
（请想想为何单核和多核的CPU效果一样。也请想想为何和cudaMemcpy的等待方式也无关）。
但是这个时间您可以用手表测出，但是您却不能用clock()测试，因为clock()测试的实际上是(*NIX下)，一个进程中的所有线程在所有的CPU核心上的占用率在时间上的积分（甚至有的时间还包含子进程的该积分值）。
所以他完全不能用来测试您的实际运行的总时间的。
（在windows上可以，因为2个实现不同)
（您可以在*nix上使用gettimeofday()来测试，以满足您的原始意图）
这个上文已经说过了，请重视一下。
感谢您的周末来访。
简要的概括说下，
一般情况下，
（1）如果原来thread1和thread2分别需要a毫秒和b毫秒，
（2）在现在的CPU上（一般双核或者多核）（这将保证无论2个线程是否100%核心占用，都将同时运行）
（3）同时thread1和thread2毫无关系，
那么您将用gettimeofday()测试出原来的a+b毫秒的运行时间变成了max(a,b)毫秒。（几乎100%概率的）
(请再次注意不要用clock(), 一些书真害人，乱写)
以及，如果它们有一定关系（例如都卡在读写磁盘上），那么此时将受限于磁盘读写瓶颈，不一定能展现加速关系了。这个也请注意。
以及，如果不需要很精确的时间关系，有个C库函数叫time(),
这个函数较为通用一点，您可以直接用:
int start = (int)time(NULL);
int stop = (int)time(NULL);
然后int elapsed = stop - start; 来获取秒数。
(你需要#include <time.h>,
请不要#include <sys/time.h>, 以本文为准。谢谢
）
（以及，您还可以用int64_t start = (int64_t)time(NULL); （需要#include <stdint.h>和<time.h>)
以可以在更遥远的未来运行您的代码(2038年以后, 短期用int足够)
）
此函数简单易用，对秒级的计时，推荐用这个。
更精确的计时，在贵平台上，请使用gettimeofday()。谢谢。
谢谢您的提醒，确实clock是个问题啊。目前正在熟悉那个Cuda Tool的使用，是个强大的工具，呵呵！
我还有一个问题：现在是两个线程在CPU端运行，一个线程使用GPU运行核函数，另外一个线程就不能同时使用GPU运行核函数了，有没有一种方法，当第一个线程使用GPU完后，第二个线程马上就能知道然后继续使用呢？
您可以2个host thread使用2个流的（参考：手册中的流部分，例如cudaStreamCreate()等)，
来<<<>>>2个kernel的：
如果:
(1)这2个kernel都足够小，那么它们将同时运行。
(2)反之，在一个kernel结束的时候，另外一个kernel将自动立刻开始自行。
您可以使用这个方法。