用户
 找回密码
 立即注册
bbskate1990 该用户已被删除
发表于 2013-11-19 10:02:39
2283115
最近导师弄个个kelper的架构的显卡给我们弄 主要想尝试一下多cpu线程同时访问gpu 但是试着用openmp来开线程同时访问gpu 通过gpu observer 来观察 gpu的利用率一直没有变还是跟单线程执行相同的函数利用是一样的 执行的时间也翻倍了 也就是还是 串行来做的  想请教大大们一下有没有这方面的例程 可以借鉴参考一下啊
使用道具 举报 回复
发表于 2013-11-20 11:46:06
bbskate1990 发表于 2013-11-19 18:41
你看我这种思路 可以实现 多个线程同时访问gpu么  直接把sdk中向量相加的例子 改的 ...

LZ您好:

不能为您指出OPENMP的相关用法是否正确的建议。

以及,多个host端线程向同一个GPU发布kernel等是可以的。
但是一般这样是为了多个逻辑依赖链之间方便隔离才这样的,以及需要使用多个stream。

您的写法所有的kernel都在同一个default stream中,用的也都是同步的copy函数,这无法实现不同的stream之间相互掩盖传输的作用,也无法实现多个小线程规模的kernel同时执行的作用,而仅仅是大家都往同一个GPU上发布任务而已。

大致如此,祝您好运~
使用道具 举报 回复 支持 反对
发表于 2013-11-19 10:26:41
LZ您好:

请您提供您的具体实现的写法,以便分析。

另外,请不要使用默认stream,默认stream中隐含的同步机制会破坏多stream的并行性。

祝您好运~
使用道具 举报 回复 支持 反对
发表于 2013-11-19 18:41:01
ice 发表于 2013-11-19 10:26
LZ您好:

请您提供您的具体实现的写法,以便分析。
  1. #include<stdio.h>
  2. #include<cuda_runtime.h>
  3. #include<omp.h>
  4. #include<time.h>
  5. #include<opencv2/core/core.hpp>
  6. #include<iostream>

  7. __global__ void
  8. vectorAdd(const float *A,const float *B, float *C,int numElements)
  9. {
  10.         int i = blockDim.x * blockIdx.x + threadIdx.x;
  11.         if (i < numElements)
  12.         {
  13.                  C[i] = A[i] * B[i];
  14.         }
  15. }

  16. int main(void)
  17. {
  18.         
  19.         cudaError_t err = cudaSuccess;
  20.         int numElements = 5000000;
  21.         size_t size = numElements * sizeof(float);
  22.         printf("[Vector addition of %d elements]\n", numElements);

  23.         float *h_a = (float *)malloc(size);
  24.         float *h_b = (float *)malloc(size);
  25.         float *h_c = (float *)malloc(size);

  26.         for(int i = 0; i < numElements;++i)
  27.         {
  28.                 h_a[i] = rand()/(float)RAND_MAX;
  29.                 h_b[i] = rand()/(float)RAND_MAX;
  30.         }

  31.         float *d_a = NULL;
  32.         err = cudaMalloc((void**)&d_a,size);

  33.         float *d_b = NULL;
  34.         err = cudaMalloc((void **)&d_b,size);
  35.    
  36.         float *d_c = NULL;
  37.         err = cudaMalloc((void**)&d_c,size);

  38.         printf("Copy input data from the host memory to the CUDA device\n");
  39.         err = cudaMemcpy(d_a,h_a,size,cudaMemcpyHostToDevice);
  40.         err = cudaMemcpy(d_b,h_b,size,cudaMemcpyHostToDevice);
  41. /////////////////////////////////////////////////////////
  42.     float *h_a1 = (float *)malloc(size);
  43.         float *h_b1 = (float *)malloc(size);
  44.         float *h_c1 = (float *)malloc(size);

  45.         for(int i = 0; i < numElements;++i)
  46.         {
  47.                 h_a1[i] = rand()/(float)RAND_MAX;
  48.                 h_b1[i] = rand()/(float)RAND_MAX;
  49.         }

  50.         float *d_a1 = NULL;
  51.         err = cudaMalloc((void**)&d_a1,size);

  52.         float *d_b1 = NULL;
  53.         err = cudaMalloc((void **)&d_b1,size);
  54.    
  55.         float *d_c1 = NULL;
  56.         err = cudaMalloc((void**)&d_c1,size);

  57.         printf("Copy input data from the host memory to the CUDA device\n");

  58.         err = cudaMemcpy(d_a1,h_a1,size,cudaMemcpyHostToDevice);
  59.         err = cudaMemcpy(d_b1,h_b1,size,cudaMemcpyHostToDevice);
  60. /////////////////////////////////////////////////////////
  61.         int threadsPerBlock = 256;
  62.         int blocksPerGrid = (numElements + threadsPerBlock - 1)/threadsPerBlock;
  63.      //long t1 = clock();
  64.         int64 t = cv::getTickCount();
  65. #pragma omp parallel sections
  66.         {
  67. #pragma omp section
  68.                 {

  69.     printf("section1 threadid = %d\n",omp_get_thread_num());
  70.         printf("cuda kernel launch with %d blocks of %d threads\n",blocksPerGrid,threadsPerBlock);

  71.         //vectorAdd<<blocksPerGrid,threadsPerBlock>>(d_a,d_b,d_c,numElements);
  72.         vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, numElements);
  73.                 }
  74. #pragma omp section
  75.                 {
  76.         printf("section2 threadid = %d\n",omp_get_thread_num());
  77.         printf("cuda kernel launch with %d blocks of %d threads\n",blocksPerGrid,threadsPerBlock);

  78.         //vectorAdd<<blocksPerGrid,threadsPerBlock>>(d_a,d_b,d_c,numElements);
  79.         vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a1, d_b1, d_c1, numElements);
  80.                 }
  81.         }
  82.         //long t2 = clock();
  83.         //printf("cost time %ld\n",t2-t1);
  84.         std::cout<<((cv::getTickCount() - t) / cv::getTickFrequency())*1000 << " msec"<<std::endl;

  85.         err = cudaGetLastError();

  86.             if (err != cudaSuccess)
  87.     {
  88.         fprintf(stderr, "Failed to launch vectorAdd kernel (error code %s)!\n", cudaGetErrorString(err));
  89.         exit(EXIT_FAILURE);
  90.     }

  91.     // Copy the device result vector in device memory to the host result vector
  92.     // in host memory.
  93.     printf("Copy output data from the CUDA device to the host memory\n");
  94.         

  95.         err =cudaMemcpy(h_c,d_c,size,cudaMemcpyDeviceToHost);
  96.     if (err != cudaSuccess)
  97.     {
  98.         fprintf(stderr, "Failed to copy vector C from device to host (error code %s)!\n", cudaGetErrorString(err));
  99.         exit(EXIT_FAILURE);
  100.     }

  101.     // Verify that the result vector is correct
  102.     for (int i = 0; i < numElements; ++i)
  103.     {
  104.         if (fabs(h_a[i] *h_b[i] - h_c[i]) > 1e-5)
  105.         {
  106.             fprintf(stderr, "Result verification failed at element %d!\n", i);
  107.             exit(EXIT_FAILURE);
  108.         }
  109.     }
  110.     printf("Test PASSED\n");
  111.   err = cudaFree(d_a);

  112.     if (err != cudaSuccess)
  113.     {
  114.         fprintf(stderr, "Failed to free device vector A (error code %s)!\n", cudaGetErrorString(err));
  115.         exit(EXIT_FAILURE);
  116.     }
  117.     err = cudaFree(d_b);

  118.     if (err != cudaSuccess)
  119.     {
  120.         fprintf(stderr, "Failed to free device vector B (error code %s)!\n", cudaGetErrorString(err));
  121.         exit(EXIT_FAILURE);
  122.     }
  123.     err = cudaFree(d_c);

  124.     if (err != cudaSuccess)
  125.     {
  126.         fprintf(stderr, "Failed to free device vector C (error code %s)!\n", cudaGetErrorString(err));
  127.         exit(EXIT_FAILURE);
  128.     }

  129.     // Free host memory
  130.     free(h_a);
  131.     free(h_b);
  132.     free(h_c);

  133.     // Reset the device and exit
  134.     err = cudaDeviceReset();

  135.     if (err != cudaSuccess)
  136.     {
  137.         fprintf(stderr, "Failed to deinitialize the device! error=%s\n", cudaGetErrorString(err));
  138.         exit(EXIT_FAILURE);
  139.     }

  140.     printf("Done\n");
  141.     return 0;
  142. }
复制代码
你看我这种思路 可以实现 多个线程同时访问gpu么  直接把sdk中向量相加的例子 改的
使用道具 举报 回复 支持 反对
发表于 2013-11-20 13:53:43
ice 发表于 2013-11-20 11:46
LZ您好:

不能为您指出OPENMP的相关用法是否正确的建议。

太感谢您了 我大概明白要在哪方面 下功夫  
使用道具 举报 回复 支持 反对
发表于 2013-11-20 13:55:37
bbskate1990 发表于 2013-11-20 13:53
太感谢您了 我大概明白要在哪方面 下功夫

不客气的,欢迎您常来~
使用道具 举报 回复 支持 反对
发表于 2013-11-25 15:48:54

哥们儿 我按照您的思路写了一段代码 但在nsight中 观察到stre

本帖最后由 bbskate1990 于 2013-11-25 15:53 编辑
ice 发表于 2013-11-20 13:55
不客气的,欢迎您常来~
  1. #include<cuda.h>
  2. #include<stdio.h>
  3. #include<stdlib.h>
  4. #include<cv.h>
  5. #include<highgui.h>
  6. #include<cxcore.h>
  7. #include<math_functions.h>
  8. #include<cuda_runtime.h>
  9. #include<time.h>

  10. #include<iostream>


  11. const char *filename[16] = {"test1.bmp","test2.bmp","test3.bmp","test4.bmp","test5.bmp","test6.bmp","test7.bmp","test8.bmp","test9.bmp","test10.bmp","test11.bmp","test12.bmp","test13.bmp","test14.bmp","test15.bmp","test16.bmp"};
  12. const int nstreams = 16;
  13. const int interation = 20;
  14. void checkCUDAError(const char* msg);

  15. __global__ void edge_gpu(unsigned char* buff , unsigned char* buffer_out , int w , int h)
  16. {
  17.         int x = blockIdx.x * blockDim.x +threadIdx.x ;
  18.         int y = blockIdx.y * blockDim.y +threadIdx.y;
  19.         int width = w , height = h ;

  20.         if((x>=0 && x < width) && (y>=0 && y<height))
  21.         {
  22.                 int hx = -buff[width*(y-1) + (x-1)] + buff[width*(y-1)+(x+1)]
  23.                 -2*buff[width*(y)+(x-1)] + 2*buff[width*(y)+(x+1)]
  24.                 -buff[width*(y+1)+(x-1)] + buff[width*(y+1)+(x+1)];

  25.                 int vx = buff[width*(y-1)+(x-1)] +2*buff[width*(y-1)+(x+1)] +buff[width*(y-1)+(x+1)]
  26.                 -buff[width*(y+1)+(x-1)] - 2*buff[width*(y+1)+(x)] - buff[width*(y+1)+(x+1)];


  27.                 int val = (int)sqrt((float)(hx) * (float)(hx) + (float)(vx) * (float)(vx));                                       

  28.                 buffer_out[y * width + x] = (unsigned char) val;                                                        
  29.         }
  30. }


  31. void checkCUDAError(const char* msg)
  32. {
  33.         cudaError_t err = cudaGetLastError();
  34.         if (cudaSuccess != err)
  35.         {
  36.                 fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString(err));
  37.                 exit(EXIT_FAILURE);
  38.         }
  39. }



  40. int main(int argc , char** argv)
  41. {
  42.         int nWidth = 2448;
  43.         int nHeight = 2048;
  44.         IplImage *temp[nstreams];
  45.         for(int i = 0;i < nstreams;i++)
  46.                 temp[i] = cvLoadImage(filename[i],0);

  47.         int buffer_size = nWidth * nHeight;

  48.         unsigned char *buffer[nstreams];
  49.         unsigned char* buf[nstreams] ;
  50.         unsigned char *buffer_dev[nstreams];
  51.         unsigned char *buffer_out[nstreams];

  52.         for(int i = 0;i < nstreams;i++)
  53.                 buffer[i] = (unsigned char*)(temp[i]->imageData);


  54.         for(int i = 0;i < nstreams;i++){
  55.                 //buf[i] = (unsigned char*) malloc(buffer_size);
  56.                 cudaMallocHost((void**)&buf[i],buffer_size);
  57.                 checkCUDAError("Memory Allocation");
  58.         }
  59.         for(int i = 0;i < nstreams;i++){
  60.                 cudaMalloc((void**)&buffer_out[i],buffer_size);
  61.                 checkCUDAError("Memory Allocation");
  62.         }

  63.         for(int i = 0;i < nstreams;i++){
  64.                 cudaMalloc((void**)&buffer_dev[i],buffer_size);
  65.                 checkCUDAError("Memory Allocation");
  66.         }




  67.         //==========  create a stream  ==========

  68.         cudaStream_t *streams = (cudaStream_t*)malloc(nstreams * sizeof(cudaStream_t));
  69.         for(int i = 0;i < nstreams;i++)
  70.                 cudaStreamCreate(&(streams[i]));
  71.         //=========================================        

  72.         for(int i = 0;i< nstreams;i++){
  73.                 cudaMemcpy(buffer_dev[i] , buffer[i] , buffer_size , cudaMemcpyHostToDevice);
  74.                 std::cout<<i<<std::endl;
  75.                 checkCUDAError("Memory Copy From Host To Device");
  76.         }
  77.         dim3 threadsPerBlock(8,8);
  78.         dim3 numBlocks((nWidth)/8,(nHeight)/8);

  79.         for(int i = 0;i < nstreams;i++){
  80.                 edge_gpu<<< numBlocks , threadsPerBlock , 0 , streams[i] >>>(buffer_dev[i] , buffer_out[i], nWidth , nHeight);
  81.                 checkCUDAError("Kernel");
  82.         }

  83.         //for (int i = 0; i < nstreams; i ++)
  84.         //        cudaStreamSynchronize(streams[i]);

  85.         for(int i = 0;i < nstreams;i++){
  86.                 cudaMemcpy(buf[i] , buffer_out[i] , buffer_size  , cudaMemcpyDeviceToHost);
  87.                 checkCUDAError("Memory Copy From Device To Host");
  88.         }


  89.         //==========Check the result==============
  90.         //memcpy(buffer[1],buf[1],buffer_size);//copying memory to show image
  91.         //cvNamedWindow("0");
  92.         //cvShowImage("0",temp[1]);
  93.         //cvWaitKey(0);
  94.         for(int i = 0;i < nstreams;i++)
  95.                 cudaFree(&(buffer_dev[i]));

  96.         for(int i = 0;i < nstreams;i++)
  97.                 cudaFree(&(buffer_out[i]));

  98.         for(int i = 0;i < nstreams;i++)
  99.                 cudaStreamDestroy(streams[i]);        

  100.         //for(int i = 0;i < nstreams;i++)
  101.         //free(&(buf[i]));
  102.         for(int i = 0;i < nstreams;i++)
  103.                 cudaFreeHost(buf[i]);

  104.         for(int i = 0;i < nstreams;i++)
  105.                 cvReleaseImage(&(temp[i]));

  106.         free(streams);

  107.         printf("\n\nDONE...!!...Copy Successful...!!\n\n");

  108.         return 1;
  109. }
复制代码

本帖子中包含更多资源

您需要 登录 才可以下载或查看,没有帐号?立即注册

x
使用道具 举报 回复 支持 反对
发表于 2013-11-25 15:53:22
楼主您好,

有问题请直接询问。

不要贴个图片就走人,让人猜测您心里的问题是什么。

谢谢。
使用道具 举报 回复 支持 反对
发表于 2013-11-25 15:54:47
ice 发表于 2013-11-20 13:55
不客气的,欢迎您常来~

哥们 我按照您的思路写了一段图像处理的 小代码  但是在nsight中调试 观察到stream 并不是并行的 我用的计算能力为3.5的显卡 而且sdk中hyperq的小例子调试出来也是并行的 您能指点一下 我这段代码哪儿有问题么  谢谢了  代码具体在上一楼
使用道具 举报 回复 支持 反对
发表于 2013-11-25 15:57:04
玫瑰幻想 发表于 2013-11-25 15:53
楼主您好,

有问题请直接询问。

您好不好意思 我编辑的 有问题 所以贴在下一楼了   我想问 我用的是kepler架构的显卡 我这段代码 stream 为啥无法并行   
使用道具 举报 回复 支持 反对
12下一页
发新帖
您需要登录后才可以回帖 登录 | 立即注册