版主您好，我又测试了段代码，还有一些问题想请教：如果使用18行，结果显示为0.00；如果使用19行，结果显示为110.00； kernel中定义变量应该是没问题的吧？而且每个线程执行的时候，这些变量是不是应该相互不影响；那么是精度的问题吗？
#include <stdio.h>
#include <cuda_runtime.h>
#include "..\..\..\..\..\..\..\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\include\device_launch_parameters.h"
#include "..\..\..\..\..\..\..\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\include\common_functions.h"
#define Nx 10
#define Nz 10
__global__ void Mykernel(float *AA)
{
float r=2000.0;
float tt=0.001;
int ix=threadIdx.x+blockDim.x*blockIdx.x;
int iz=threadIdx.y+blockDim.y*blockIdx.y;
if(ix<Nx&&iz<Nz){
//AA[iz*Nx+ix] = 1.0/r*tt+110.0;
AA[iz*Nx+ix] =1.0/2000.0*0.001+110.0;
}
}
int main()
{
float *CPU_V;
cudaMallocHost((void**)&CPU_V,Nx*Nz*sizeof(float));
float *GPU_V;
cudaMalloc((void**)&GPU_V,Nx*Nz*sizeof(float)); cudaMemset(GPU_V,0.0,Nx*Nz*sizeof(float));
dim3 TheadsPerBlock(32,32);
dim3 BlockPerGrid((Nx+32-1)/32,(Nz+32-1)/32);
Mykernel<<<BlockPerGrid,TheadsPerBlock >>>(GPU_V);
cudaError_t error;
error=cudaDeviceSynchronize();
cudaError_t error_cy;
error_cy=cudaMemcpy(CPU_V,GPU_V,Nx*Nz*sizeof(float),cudaMemcpyDeviceToHost);
printf("error=%s,error_cy=%s\n",cudaGetErrorString(error),cudaGetErrorString(error_cy));
for (int iz=0;iz<Nz;iz++)
{
for (int ix=0;ix<Nx;ix++)
{
//if (CPU_V[iz*Nx+ix]!=211.0)
{
printf("%.10f,iz=%d,ix=%d\n",CPU_V[iz*Nx+ix],iz,ix);
}
}
}
cudaFree(CPU_V);
//cudaDeviceReset();
cudaFree(GPU_V);
printf("hello world!\n");
return 0;
}
LZ您好：
float精度的浮点数，换算成10进制，只有大概7位有效数字，按照您的写法，1.0/2000.0*0.001+110，需要10位有效数字才能表示，所以您使用double精度的浮点数可以正确表达，而使用float精度的浮点数则会丢失部分尾数。
另，就您18行的写法，经我简单测试，是可以在kernel里面这样计算的，并无问题（和19行一样会丢失尾数）。
大致如此，祝您编码顺利~
斑竹，您好，感谢您的回复
LZ您好：
您本楼的描述和您1#的描述是不一致的，为了避免进一步的混淆，请给出您本楼描述的（“所有都换成double”）情况下的具体的完整的代码，以便分析。
祝您好运~
版主你好，代码如下：
#include <stdio.h>
#include <cuda_runtime.h>
#include "..\..\..\..\..\..\..\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\include\device_launch_parameters.h"
#include "..\..\..\..\..\..\..\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\include\common_functions.h"
#define Nx 10
#define Nz 10
__global__ void Mykernel(double *AA)
{
double r=2000.0;
double tt=0.001;
int ix=threadIdx.x+blockDim.x*blockIdx.x;
int iz=threadIdx.y+blockDim.y*blockIdx.y;
if(ix<Nx&&iz<Nz){
AA[iz*Nx+ix] = 1.0/r*tt+110.0;
//AA[iz*Nx+ix] =1.0/2000.0*0.001+110.0;
}
}
int main()
{
double *CPU_V;
cudaMallocHost((void**)&CPU_V,Nx*Nz*sizeof(double));
double *GPU_V;
cudaMalloc((void**)&GPU_V,Nx*Nz*sizeof(double)); cudaMemset(GPU_V,0.0,Nx*Nz*sizeof(double));
dim3 TheadsPerBlock(32,32);
dim3 BlockPerGrid((Nx+32-1)/32,(Nz+32-1)/32);
Mykernel<<<BlockPerGrid,TheadsPerBlock >>>(GPU_V);
cudaError_t error;
error=cudaDeviceSynchronize();
cudaError_t error_cy;
error_cy=cudaMemcpy(CPU_V,GPU_V,Nx*Nz*sizeof(double),cudaMemcpyDeviceToHost);
printf("error=%s,error_cy=%s\n",cudaGetErrorString(error),cudaGetErrorString(error_cy));
for (int iz=0;iz<Nz;iz++)
{
for (int ix=0;ix<Nx;ix++)
{
//if (CPU_V[iz*Nx+ix]!=211.0)
{
printf("CPU_V=%.10lf,iz=%d,ix=%d\n",CPU_V[iz*Nx+ix],iz,ix);
}
}
}
cudaFree(CPU_V);
//cudaDeviceReset();
cudaFree(GPU_V);
printf("hello world!\n");
return 0;
}
刚才又试了一下，如果直接定义double r_daoshu=0.0005;
LZ您好：
GPU计算除法仅仅是慢点，不是不能计算。
我也觉的应该是这样，可是为什么5#的代码 在我电脑上跑出来就不对呢
LZ您好：
经过仔细调试您的代码，发现问题如下：
1：您的问题在于您的kernel挂掉了。通过在您的kernel后面添加一句
2：如果您将编译类型转换为“release”，可以立即得到正确的结果。
3：考虑到您实际上只计算了1010的数据，但是却分配了32 32的线程，90%以上的线程都只是去占用资源但是不干活，所以您如果使用较小的block（比如10*10的block），那么在debug下也可以正确运行。
4：考虑到您的问题出在计算除法的地方，而除法在GPU中是一个有很多中间过程的计算，需要使用较多的寄存器，而您block设置又达到了上限，可能导致每个线程没有足够的寄存器资源。一般这种情况，如果是预先可知的，请设置较小的block尺寸。
如果您必须使用较大的block尺寸，并且不在乎可能在速度方面的降低，那么也可以通过限制Max Used Register来强制GPU使用local memory，从而达到“可用”（但非“好用”）的目的。
5：以及为何在release下能直接得到正确的结果，一方面release编译可以按照节省资源的方式编译，无需保留临时信息供调试，另一方面，编译器可能直接识别您是除以一个常数，从而直接优化掉这个除法，编译过后就直接是乘以该数的倒数的形式。
您的问题完整分析如上，祝您编码顺利~
system
2013 年11 月 7 日 06:02
10
刚才和玫瑰斑竹又研究和测试了一下：
1：我之前在CUDA 5.0上编译的release版的代码编译器并未优化掉除法，而是生成了完整的除法代码。
2：您1#中的代码，如果在2000.0和0.001的地方主动指定为2000.0f和0.001f，那么可能直接解决问题。（此时5.0及更早版本的编译器将使用您指定的单精度版的数值进行计算，而不是按照双精度计算，最后截断，更省资源。）
3：建议您直接升级到CUDA 5.5，在CUDA 5.5中，经观察，上述对常量的除法将直接被编译掉，无论是debug模式还是release模式，速度更快，效果更好。对于赋值给单精度浮点变量的双精度值，也将直接就地在编译的时候转换为单精度的值。
综上，推荐您升级到CUDA 5.5，以直接解决各种困扰。
祝您编码顺利~
system
2013 年11 月 7 日 07:24
11
感谢ice版主和玫瑰斑细心测试与专业的回答，将block规模改小了后确实可以了，我以后在编程的时候会注意这一点
system
2013 年11 月 7 日 08:00
12
LZ您好：
因为实际上kernel除了一般出现的异步错误以外，还可能出现在调用时就返回的同步错误（比较少见），cudaDeviceSynchronize()只能捕获kernel启动的异步错误，而runtime API中kernel是void类型的，无返回值，因此同步错误就只能通过cudaGetLastError()来寻找了。
比如本例，您的双精度的除法被编译成了使用大量寄存器的实现，在编译的时候，编译器并不知道最终会上多少个线程，而您实际上的单一block的线程数量太多，因而在启动kernel的时候立即就会发现无法执行（对于fermi和kepler SM3.5的GPU，一个block上最大数量的线程与每个线程上最大数量的寄存器，这两个限制不能同时满足。）。而如果您手工控制单一线程使用的寄存器数量，使得一个1024线程的线程所使用的寄存器的数量不超过一个SM的寄存器资源的总量的话，kernel是可以成功启动的。
同时，因为这是一个同步Error，所以并没有像异步Error一样导致后续的cudaMemcpy()失效，这样被您cudaMemset()d GPU_V数组的值就被复制回来了。
大致分析如此，亦有玫瑰斑竹的贡献。
总而言之，您需要两者配合使用，才能对付一些比较少见的疑难情况。
欢迎您莅临cudazone@China，祝您编码愉快~
system
2013 年11 月 7 日 10:57
13
好的 ，又学到不少东西
system
2013 年11 月 7 日 11:01
14
LZ您好：
不客气的，您卸载5.0然后安装一份5.5就可以，或者您也可以直接安装CUDA 5.5，CUDA Toolkit是可以同时存在多个版本的，但是一些环境变量会被后装的那个修改，一般无问题，你要是不放心的话，只保留5.5也是很好的选择。
现在的CUDA Toolkit都是直接安装就可以免配置使用的。
大致如此，祝您好运~