简单cuda程序

system · 2013 年7 月 14 日 14:07

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
 
#include <iostream>
using namespace std;
 
#define Q 10
__global__ void kernel(float *df[Q])
{
   int i,j;
   for (i=0;i<10;i++)
   {
   for (j=0;j<10;j++)
   {
   df[i][j]*=2;
   }
   }
}
int main()
{
   float *f[Q];
   int i,j;
   for (i=0;i<Q;i++)
   {
   f[i]=(float *)malloc(sizeof(float)*Q);
   }
   for (i=0;i<10;i++)
   {
   for (j=0;j<10;j++)
   {
   f[i][j]=i+j;
   }
   }
 
   float *df[Q];
   for (i=0;i<Q;i++)
   {
   cudaMalloc((void **) &df[i],sizeof(float)*Q);
   cudaMemcpy(df[i],f[i],sizeof(float)*Q,cudaMemcpyHostToDevice);
   }
   kernel<<<1,1>>>(df);
 
   for (i=0;i<Q;i++)
   {
   cudaMemcpy(f[i],df[i],sizeof(float)*Q,cudaMemcpyDeviceToHost);
   cudaFree(df[i]);
   }
 
   cout << f[1][2] << endl;
   getchar();
   for (i=0;i<Q;i++)
   {
   free(f[i]);
   }
   return 0;
}

kernel函数的就是把f中的元素都乘以2，最后输出的f[1][2]应该是6才对，我的结果是3，不知道哪里出错了。另外，核函数这么调用是否正确kernel<<<1,1>>>(df);

system · 2013 年7 月 14 日 14:55

LZ您好：

如果您在kernel后面加上cudaDeviceSynchronize()并检查此函数的返回值的话，您会发现您的kernel没有正常运行，以及后续的cudaMemcpy也不会继续运行（这是CUDA的特色）。

您的问题在于您给kernel提供的参数是指向host端数组df的指针，这样kernel必挂。您必须提供一个指向device端存储空间的但本身保存在host端的指针作为kernel的参数才行。

大致分析如此，祝您好运~

system · 2013 年7 月 15 日 01:50

非常感谢您的回答。我大致理解了你的意思，我做过尝试，把df数组的元素单独传给kernel，也就是kernel<<<1,1>>>(df[0], df[1], … , df[9])，这样修改就可以得到正确结果。
df只是数组的首个元素（也是地址）的地址，df≠df[0]，这样导致cuda无法识别出错。
但是如果数组的元素数目过大，假如100个，也不能直接将100个元素都列出来吧？！kernel<<<1,1>>>(df[0], df[1], … , df[99])，请问版主，有其他的处理办法吗？

system · 2013 年7 月 15 日 04:57

LZ您好：

您的指针数组df是在host端内存上的，里面的各个指针元素自身是存在host端内存的，但是其保存的地址（也就是其指向）是device端的地址（因为是cudaMalloc()申请的）。

那么您的数组的每个元素，满足kernel函数参数使用指针的要求，即“自身存储于host端，但是指向（保存的地址）是device端地址”。
而如果您将df数组的地址作为参数传递给kernel，那么这个地址指向的是host端的存储空间，此时kernel必挂。

这是您代码问题的本源。

system · 2013 年7 月 15 日 05:06

以及，解决您的问题原则上至少有两种方法：

1：依然按照您的思路，您需要在device端申请一段空间来保存df的一个副本，并在host端定义一个指向指针数组的指针来指向这个df的副本，并将df的内容复制给这个副本。然后使用这个指向指针数组的指针作为kernel的参数，此时可以正常执行。

2：在kernel中使用两维数组一般无需这样定义一堆指针去手工维护，直接申请一段连续的线性缓冲区，并使用带有维度信息的指针即可。（使用普通指针，并自己计算偏移量也很容易实现，并无其他性能代价。）
下图是一个例子，供您参考：
[attach]3263[/attach]

祝您编码顺利~

system · 2013 年7 月 15 日 07:17

非常感谢版主细心又专业的回答，我按照您给提供的第一种思路，得到了正确的答案，代码如下：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
using namespace std;
#define Q 10
__global__ void kernel(float *df[Q])
{
	int i,j;
	for (i=0;i<10;i++)
	{
		for (j=0;j<10;j++)
		{
			df[i][j]*=2;
		}
	}
}
int main()
{
	float *f[Q];
	int i,j;
	for (i=0;i<Q;i++)
	{
		f[i]=(float *)malloc(sizeof(float)*Q);
	}
	for (i=0;i<10;i++)
	{
		for (j=0;j<10;j++)
		{
			f[i][j]=i+j;
		}
	}

	float *df[Q];
	for (i=0;i<Q;i++)
	{
		cudaMalloc((void **) &df[i],sizeof(float)*Q);
		cudaMemcpy(df[i],f[i],sizeof(float)*Q,cudaMemcpyHostToDevice);
	}
	float ** ddf;
	cudaMalloc((void **) &ddf, sizeof(float *)*Q);
	cudaMemcpy(ddf,df,sizeof(float *)*Q,cudaMemcpyHostToDevice);

	kernel<<<1,1>>>(ddf);

	for (i=0;i<Q;i++)
	{
		cudaMemcpy(f[i],df[i],sizeof(float)*Q,cudaMemcpyDeviceToHost);
		cudaFree(df[i]);
	}
	cudaFree(ddf);
	cout << f[1][2] << endl;
	getchar();
	for (i=0;i<Q;i++)
	{
		free(f[i]);
	}
	return 0;
}

您看看还有需要优化的吗？
另外：cuda c programing guide中有个矩阵加法的例子（guide p.18），文档中没有给出完整的代码，我想应该可以使用您提供的第二种思路来完成，right？！
[attach]3264[/attach]

system · 2013 年7 月 15 日 09:10

LZ您好：

1：您修改后的代码是正确的，以及这个只是个示例性质的代码，谈不上什么优化。以及就使用二维数组而言，一般建议第二种方法，要简单得多。

2：LZ您好，应该可以的，包括您原本的程序，也可以考虑使用并推荐使用第二种方法的。

祝您编码顺利~

system · 2013 年7 月 15 日 12:28

版主您好，我用您说的第二种思路也做出来了，并且完成了cuda guide中矩阵加法的计算，如下：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
using namespace std;
#define M 10
#define N 15

__global__ void matrixAdd(float A[M][N], float B[M][N], float C[M][N])
{
	int i=threadIdx.x;
	int j=threadIdx.y;
	C[i][j]=A[i][j]+B[i][j];
}
int main()
{
	float a[M][N],b[M][N],c[M][N];
	float *A,*B,*C;
	cudaMalloc((void **) &A, sizeof(float)*M*N);
	cudaMalloc((void **) &B, sizeof(float)*M*N);
	cudaMalloc((void **) &C, sizeof(float)*M*N);
	for (int j=0;j<M;j++)
	{
		for (int i=0;i<N;i++)
		{
			a[j][i]=i+j;
			b[j][i]=i+2*j;
		}
	}

	//cudaMemcpy(A, (float *)a, sizeof(float)*M*N, cudaMemcpyHostToDevice);
	//cudaMemcpy(B, (float *)b, sizeof(float)*M*N, cudaMemcpyHostToDevice);
	cudaMemcpy(A, a, sizeof(float)*M*N, cudaMemcpyHostToDevice);
	cudaMemcpy(B, b, sizeof(float)*M*N, cudaMemcpyHostToDevice);


	int numBlocks = 1;
	dim3 threadsPerBlock(M, N);
	matrixAdd<<<numBlocks, threadsPerBlock>>>((float (*)[N])A, (float (*)[N])B, (float (*)[N])C);


	//cudaMemcpy((float *)c, C, sizeof(float)*M*N, cudaMemcpyDeviceToHost);
	cudaMemcpy(c, C, sizeof(float)*M*N, cudaMemcpyDeviceToHost);

	cout << c[1][2] << endl;
	cudaFree(A);
	cudaFree(B);
	cudaFree(C);

	return 0;
}

我有个疑问，如何才能像书上描述的那样，核函数的调用为matrixAdd<<<numBlocks, threadsPerBlock>>>(A, B, C); 并且保证核函数中可以使用二维数组？

system · 2013 年7 月 15 日 12:36

LZ您好：

看到您新给出的实现了，这个实现是正确的。

以及，我在给您那个示例里面，实际上是先定义了一个一般的不含维度信息的指针，然后在调用kernel的时候对参数就地进行了类型转换，加入了维度信息。

您也可以直接定义带有维度信息的指针，这样在作为kernel参数的时候，就不用再次转换了。

即：
您当前的写法是：定义为 float *A；

调用的时候就地转换类型为：(float(*)[N])A

您可以更改为：

定义指针为：float (*A) [N];

调用的时候，直接使用A作为参数，kernel声明的时候参数类型还保持原来写法不变。

这样就可以在调用kenrel函数的时候，直接使用指针名而不用另外进行转换了。

大致如此，供您参考。

祝您编码顺利~

system · 2013 年7 月 15 日 13:45

版主您好，我改了一下，确实如此。
这个贴里面的代码虽然很简单，但是通过跟您的交流，我也获益匪浅，再次感谢！

system · 2013 年7 月 15 日 16:01

LZ您好：

恭喜您进一步弄清楚了这个小问题，积跬步以至千里，汇细流以成江海，祝您早日达到更高的境界！

以及十分荣幸为您服务，欢迎您常来论坛。

祝您好运~