最近导师弄个个kelper的架构的显卡给我们弄 主要想尝试一下多cpu线程同时访问gpu 但是试着用openmp来开线程同时访问gpu 通过gpu observer 来观察 gpu的利用率一直没有变还是跟单线程执行相同的函数利用是一样的 执行的时间也翻倍了 也就是还是 串行来做的 想请教大大们一下有没有这方面的例程 可以借鉴参考一下啊
LZ您好：
请您提供您的具体实现的写法，以便分析。
另外，请不要使用默认stream，默认stream中隐含的同步机制会破坏多stream的并行性。
祝您好运~
#include<stdio.h>
#include<cuda_runtime.h>
#include<omp.h>
#include<time.h>
#include<opencv2/core/core.hpp>
#include<iostream>
__global__ void
vectorAdd(const float *A,const float *B, float *C,int numElements)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < numElements)
{
C[i] = A[i] * B[i];
}
}
int main(void)
{
cudaError_t err = cudaSuccess;
int numElements = 5000000;
size_t size = numElements * sizeof(float);
printf("[Vector addition of %d elements]\n", numElements);
float *h_a = (float *)malloc(size);
float *h_b = (float *)malloc(size);
float *h_c = (float *)malloc(size);
for(int i = 0; i < numElements;++i)
{
h_a[i] = rand()/(float)RAND_MAX;
h_b[i] = rand()/(float)RAND_MAX;
}
float *d_a = NULL;
err = cudaMalloc((void**)&d_a,size);
float *d_b = NULL;
err = cudaMalloc((void **)&d_b,size);
float *d_c = NULL;
err = cudaMalloc((void**)&d_c,size);
printf("Copy input data from the host memory to the CUDA device\n");
err = cudaMemcpy(d_a,h_a,size,cudaMemcpyHostToDevice);
err = cudaMemcpy(d_b,h_b,size,cudaMemcpyHostToDevice);
/////////////////////////////////////////////////////////
float *h_a1 = (float *)malloc(size);
float *h_b1 = (float *)malloc(size);
float *h_c1 = (float *)malloc(size);
for(int i = 0; i < numElements;++i)
{
h_a1[i] = rand()/(float)RAND_MAX;
h_b1[i] = rand()/(float)RAND_MAX;
}
float *d_a1 = NULL;
err = cudaMalloc((void**)&d_a1,size);
float *d_b1 = NULL;
err = cudaMalloc((void **)&d_b1,size);
float *d_c1 = NULL;
err = cudaMalloc((void**)&d_c1,size);
printf("Copy input data from the host memory to the CUDA device\n");
err = cudaMemcpy(d_a1,h_a1,size,cudaMemcpyHostToDevice);
err = cudaMemcpy(d_b1,h_b1,size,cudaMemcpyHostToDevice);
/////////////////////////////////////////////////////////
int threadsPerBlock = 256;
int blocksPerGrid = (numElements + threadsPerBlock - 1)/threadsPerBlock;
//long t1 = clock();
int64 t = cv::getTickCount();
#pragma omp parallel sections
{
#pragma omp section
{
printf("section1 threadid = %d\n",omp_get_thread_num());
printf("cuda kernel launch with %d blocks of %d threads\n",blocksPerGrid,threadsPerBlock);
//vectorAdd<<blocksPerGrid,threadsPerBlock>>(d_a,d_b,d_c,numElements);
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, numElements);
}
#pragma omp section
{
printf("section2 threadid = %d\n",omp_get_thread_num());
printf("cuda kernel launch with %d blocks of %d threads\n",blocksPerGrid,threadsPerBlock);
//vectorAdd<<blocksPerGrid,threadsPerBlock>>(d_a,d_b,d_c,numElements);
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a1, d_b1, d_c1, numElements);
}
}
//long t2 = clock();
//printf("cost time %ld\n",t2-t1);
std::cout<<((cv::getTickCount() - t) / cv::getTickFrequency())*1000 << " msec"<<std::endl;
err = cudaGetLastError();
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to launch vectorAdd kernel (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
// Copy the device result vector in device memory to the host result vector
// in host memory.
printf("Copy output data from the CUDA device to the host memory\n");
err =cudaMemcpy(h_c,d_c,size,cudaMemcpyDeviceToHost);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to copy vector C from device to host (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
// Verify that the result vector is correct
for (int i = 0; i < numElements; ++i)
{
if (fabs(h_a[i] *h_b[i] - h_c[i]) > 1e-5)
{
fprintf(stderr, "Result verification failed at element %d!\n", i);
exit(EXIT_FAILURE);
}
}
printf("Test PASSED\n");
err = cudaFree(d_a);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to free device vector A (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
err = cudaFree(d_b);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to free device vector B (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
err = cudaFree(d_c);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to free device vector C (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
// Free host memory
free(h_a);
free(h_b);
free(h_c);
// Reset the device and exit
err = cudaDeviceReset();
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to deinitialize the device! error=%s\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
printf("Done\n");
return 0;
}
你看我这种思路 可以实现 多个线程同时访问gpu么 直接把sdk中向量相加的例子 改的
LZ您好：
不能为您指出OPENMP的相关用法是否正确的建议。
以及，多个host端线程向同一个GPU发布kernel等是可以的。
但是一般这样是为了多个逻辑依赖链之间方便隔离才这样的，以及需要使用多个stream。
您的写法所有的kernel都在同一个default stream中，用的也都是同步的copy函数，这无法实现不同的stream之间相互掩盖传输的作用，也无法实现多个小线程规模的kernel同时执行的作用，而仅仅是大家都往同一个GPU上发布任务而已。
大致如此，祝您好运~
太感谢您了 我大概明白要在哪方面 下功夫
不客气的，欢迎您常来~
#include<cuda.h>
#include<stdio.h>
#include<stdlib.h>
#include<cv.h>
#include<highgui.h>
#include<cxcore.h>
#include<math_functions.h>
#include<cuda_runtime.h>
#include<time.h>
#include<iostream>
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"};
const int nstreams = 16;
const int interation = 20;
void checkCUDAError(const char* msg);
__global__ void edge_gpu(unsigned char* buff , unsigned char* buffer_out , int w , int h)
{
int x = blockIdx.x * blockDim.x +threadIdx.x ;
int y = blockIdx.y * blockDim.y +threadIdx.y;
int width = w , height = h ;
if((x>=0 && x < width) && (y>=0 && y<height))
{
int hx = -buff[width*(y-1) + (x-1)] + buff[width*(y-1)+(x+1)]
-2*buff[width*(y)+(x-1)] + 2*buff[width*(y)+(x+1)]
-buff[width*(y+1)+(x-1)] + buff[width*(y+1)+(x+1)];
int vx = buff[width*(y-1)+(x-1)] +2*buff[width*(y-1)+(x+1)] +buff[width*(y-1)+(x+1)]
-buff[width*(y+1)+(x-1)] - 2*buff[width*(y+1)+(x)] - buff[width*(y+1)+(x+1)];
int val = (int)sqrt((float)(hx) * (float)(hx) + (float)(vx) * (float)(vx));
buffer_out[y * width + x] = (unsigned char) val;
}
}
void checkCUDAError(const char* msg)
{
cudaError_t err = cudaGetLastError();
if (cudaSuccess != err)
{
fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
}
int main(int argc , char** argv)
{
int nWidth = 2448;
int nHeight = 2048;
IplImage *temp[nstreams];
for(int i = 0;i < nstreams;i++)
temp[i] = cvLoadImage(filename[i],0);
int buffer_size = nWidth * nHeight;
unsigned char *buffer[nstreams];
unsigned char* buf[nstreams] ;
unsigned char *buffer_dev[nstreams];
unsigned char *buffer_out[nstreams];
for(int i = 0;i < nstreams;i++)
buffer[i] = (unsigned char*)(temp[i]->imageData);
for(int i = 0;i < nstreams;i++){
//buf[i] = (unsigned char*) malloc(buffer_size);
cudaMallocHost((void**)&buf[i],buffer_size);
checkCUDAError("Memory Allocation");
}
for(int i = 0;i < nstreams;i++){
cudaMalloc((void**)&buffer_out[i],buffer_size);
checkCUDAError("Memory Allocation");
}
for(int i = 0;i < nstreams;i++){
cudaMalloc((void**)&buffer_dev[i],buffer_size);
checkCUDAError("Memory Allocation");
}
//========== create a stream ==========
cudaStream_t *streams = (cudaStream_t*)malloc(nstreams * sizeof(cudaStream_t));
for(int i = 0;i < nstreams;i++)
cudaStreamCreate(&(streams[i]));
//=========================================
for(int i = 0;i< nstreams;i++){
cudaMemcpy(buffer_dev[i] , buffer[i] , buffer_size , cudaMemcpyHostToDevice);
std::cout<<i<<std::endl;
checkCUDAError("Memory Copy From Host To Device");
}
dim3 threadsPerBlock(8,8);
dim3 numBlocks((nWidth)/8,(nHeight)/8);
for(int i = 0;i < nstreams;i++){
edge_gpu<<< numBlocks , threadsPerBlock , 0 , streams[i] >>>(buffer_dev[i] , buffer_out[i], nWidth , nHeight);
checkCUDAError("Kernel");
}
//for (int i = 0; i < nstreams; i ++)
// cudaStreamSynchronize(streams[i]);
for(int i = 0;i < nstreams;i++){
cudaMemcpy(buf[i] , buffer_out[i] , buffer_size , cudaMemcpyDeviceToHost);
checkCUDAError("Memory Copy From Device To Host");
}
//==========Check the result==============
//memcpy(buffer[1],buf[1],buffer_size);//copying memory to show image
//cvNamedWindow("0");
//cvShowImage("0",temp[1]);
//cvWaitKey(0);
for(int i = 0;i < nstreams;i++)
cudaFree(&(buffer_dev[i]));
for(int i = 0;i < nstreams;i++)
cudaFree(&(buffer_out[i]));
for(int i = 0;i < nstreams;i++)
cudaStreamDestroy(streams[i]);
//for(int i = 0;i < nstreams;i++)
//free(&(buf[i]));
for(int i = 0;i < nstreams;i++)
cudaFreeHost(buf[i]);
for(int i = 0;i < nstreams;i++)
cvReleaseImage(&(temp[i]));
free(streams);
printf("\n\nDONE...!!...Copy Successful...!!\n\n");
return 1;
}
楼主您好，
有问题请直接询问。
不要贴个图片就走人，让人猜测您心里的问题是什么。
谢谢。
哥们 我按照您的思路写了一段图像处理的 小代码 但是在nsight中调试 观察到stream 并不是并行的 我用的计算能力为3.5的显卡 而且sdk中hyperq的小例子调试出来也是并行的 您能指点一下 我这段代码哪儿有问题么 谢谢了 代码具体在上一楼
您好不好意思 我编辑的 有问题 所以贴在下一楼了 我想问 我用的是kepler架构的显卡 我这段代码 stream 为啥无法并行
楼主您这里引入了隐式的同步（cudaMemcpy会导致自动的同步的），所以您看不到并行非常正常。
您应当修正为使用异步版本的cudaMemcpyAsync + pinned memory, 才有可能观察到传输和复制并行。
以及，后者就算您正确的写了，您可能需要Tesla卡才能看到效果。特别从CUDA 5.0的配套驱动（及以后版本）开始，对于geforce卡的多流并行效果减弱了，更加不容易看到。
请您知悉正确的做法（第二段的cudaMemcpyAsync）， 以及可能的结果。
感谢来访。
好的 您好哥们 我刚刚开始弄 我正在看异步方面的知识 我用的是 gtx780 也可能无法看到效果么
关于是否GTX780上能看到效果，这个不一定的。
（没说普通的Geforce卡上一定看不到效果，只是说不容易看到）
空说无益，请您尝试我之前的建议(cudaMemcpyAsync + pinned memory)，并反馈结果。
好的 我做一下 在跟您讨论一下结果 thanks 
但是哥们我用sdk的example中的hyperQ的例子 在nsight中看到的 确实是并行的 感觉 那个例子 其实跟这段程序并没有大的区别啊
楼主您好，我手头没有3.5的卡，无法为您评估此Hyper-Q的效果。
我建议您自己测试下，然后发到论坛，这样可以受益很多人（包括没卡的），
鼓励他们积极买卡，共同构建CUDA的绿色生态环境。
您说呢？
感谢您的莅临。
嗯 好的 谢了