用户
 找回密码
 立即注册
afterglow 该用户已被删除
发表于 2013-11-20 17:59:07
2272311
尝试练习用driver的api,cuModuleLoadDataEx返回CUDA_ERROR_NO_BINARY_FOR_GPU
通过注销一些代码,最后发现是下面这种代码引起的
asm ("vabsdiff4.s32.u32.u32.add %0, %1, %2, %3;":"=r"(sum):"r"(a),"r"(b),"r"(previous_sum));
(注掉这种代码,cuModuleLoadDataEx可以执行成功)

问一下这种情况怎么处理比较合适?
使用道具 举报 回复
发表于 2013-11-20 22:48:24
楼主您好,

你犯了个错误:
module可以是ptx文件,也可以是cubin文件,但您无法直接载入.cu文件的。

您的.cu文件需要单独的手工通过nvcc -arch compute_xx -ptx -o your.ptx your.cu
来从您的.cu文件中生成.ptx文件的。

直接载入.cu文件(那怕只包含CUDA C写的kernel),是不可以的。
(它需要单独的nvcc离线编译一次)

感谢您的来访。
使用道具 举报 回复 支持 反对
发表于 2013-11-20 22:49:55
玫瑰幻想 发表于 2013-11-20 22:48
楼主您好,

你犯了个错误:

补充下,如果您已经是.ptx文件的,

则您应当直接写vabsdiff4....., 而无需添加asm("vabsdiff4....");
(asm只能在您的cuda c里使用,而这个是不能直接加载的,前文说过了。)
(ptx里请直接写语句即可)

感谢来访。
使用道具 举报 回复 支持 反对
发表于 2013-11-21 09:45:12
玫瑰幻想 发表于 2013-11-20 22:49
补充下,如果您已经是.ptx文件的,

则您应当直接写vabsdiff4....., 而无需添加asm("vabsdiff4....");

感谢玫瑰斑竹的热心回答

我补充一下问题的情况,我是用vs2010,cuda5.0 卡gtx780
module加载的是ptx文件,由vs2010生成的,照着例子里弄的,输出中生成ptx的指令参数有
nvcc.exe  --use-local-env --cl-version 2010 -ccbin xx  -Ixx   --keep-dir "Release" -maxrregcount=64  --machine 32 -ptx  -o my.ptx my.cu
生成的ptx文件里asm里的指令翻译成类似的代码vabsdiff4.s32.u32.u32.add %r312, %r308, %r310, %r311;

我这应该不算直接载入cu文件吧?
使用道具 举报 回复 支持 反对
发表于 2013-11-21 10:12:03
玫瑰幻想 发表于 2013-11-20 22:48
楼主您好,

你犯了个错误:

我知道了,我在cmd里单独编译ptx文件,用vs2010给的参数编译结果一致,还是返回CUDA_ERROR_NO_BINARY_FOR_GPU,您给的参数里还有-arch,加上-arch compute_35能正确加载。vs2010+cuda不是很智能。。

谢谢
使用道具 举报 回复 支持 反对
发表于 2013-11-21 11:28:56
afterglow 发表于 2013-11-21 10:12
我知道了,我在cmd里单独编译ptx文件,用vs2010给的参数编译结果一致,还是返回CUDA_ERROR_NO_BINARY_FOR ...

嗯嗯。首先恭喜楼主成功。

其次,通过您的描述,您的加载失败实际上是因为vabsdiff4在默认的1.0计算能力下不支持,导致编译失败。而编译失败,无ptx,无法加载(话说。。无ptx您加载的是什么东西?)

自然你去掉vabsdiff4或者指定了正确的计算能力后就好了。

(你1#话给我们造成了误导)
(实际原因在此)

感谢来访。
感谢来访。
使用道具 举报 回复 支持 反对
发表于 2013-11-21 12:06:08
横扫千军 发表于 2013-11-21 11:28
嗯嗯。首先恭喜楼主成功。

其次,通过您的描述,您的加载失败实际上是因为vabsdiff4在默认的1.0计算能力 ...

vs2010编译提示成功,而且确实生成了ptx文件~~
ptx文件加载很慢的样子,刚才试了试生成cubin文件,也能达到效果
最后的时间统计,runtime api和driver api的时间基本上一致
使用道具 举报 回复 支持 反对
发表于 2013-11-21 12:13:15
afterglow 发表于 2013-11-21 12:06
vs2010编译提示成功,而且确实生成了ptx文件~~
ptx文件加载很慢的样子,刚才试了试生成cubin文件,也能达 ...

楼主还编译成功,得到ptx?

那无可能的。

vabsdiff4需要至少3.0+的架构的,楼主您居然在默认的计算能力1.0上编译成功!
这违背了NV的手册,也违背了我的基本人生观和价值观。

所以您在撒谎。
但您有言论的自由,为了表示对您的自由言论的尊重,我将继续保留此贴。

但您的所有后续发帖我将不再回复。



使用道具 举报 回复 支持 反对
发表于 2013-11-21 12:37:09
玫瑰幻想 发表于 2013-11-21 12:13
楼主还编译成功,得到ptx?

那无可能的。

我确实没有说谎,为了确定你说的结论,我再一次编译(没写-arch 或者补上 -arch compute_10)确实生成了ptx文件

cu里用到了下面三个函数
__device__ int absdiff4(unsigned int a, unsigned int b, int previous_sum)
{
        int sum;
        asm ("vabsdiff4.s32.u32.u32.add %0, %1, %2, %3;":"=r"(sum):"r"(a),"r"(b),"r"(previous_sum));
        //sum += abs((int)(a    &0xff) - (int)(b    &0xff));
        //sum += abs((int)(a>>8 &0xff) - (int)(b>>8 &0xff));
        //sum += abs((int)(a>>16&0xff) - (int)(b>>16&0xff));
        //sum += abs((int)(a>>24&0xff) - (int)(b>>24&0xff));
        return sum;
}

__device__ void abs_avg(unsigned int a, unsigned int b, unsigned int c, int &sum)
{       
        asm volatile ("vavrg4.u32.u32.u32 %0, %0, %1, %0;" : "+r"(b) : "r"(c));
        asm volatile ("vabsdiff4.s32.u32.u32.add %0, %1, %2, %0;" : "+r"(sum) : "r"(a), "r"(b));
    //sum += abs( (((int)(c    &0xff) + (int)(b    &0xff) + 1) >>1) - (int)(a    &0xff));
    //sum += abs( (((int)(c>>8 &0xff) + (int)(b>>8 &0xff) + 1) >>1) - (int)(a>>8 &0xff));
    //sum += abs( (((int)(c>>16&0xff) + (int)(b>>16&0xff) + 1) >>1) - (int)(a>>16&0xff));
    //sum += abs( (((int)(c>>24&0xff) + (int)(b>>24&0xff) + 1) >>1) - (int)(a>>24&0xff));
}

__device__ void avg(unsigned int &b, unsigned int c ) // b = (b+c)/2 by byte.
{       
        asm volatile ("vavrg4.u32.u32.u32 %0, %0, %1, %0;" : "+r"(b) : "r"(c));       
    //b = ((uint32_t)(b&0xff)+(uint32_t)(c&0xff)+1) >> 1              |
    //    ((uint32_t)(b>>8 &0xff)+(uint32_t)(c>>8 &0xff)+1) >> 1 <<8  |
    //    ((uint32_t)(b>>16&0xff)+(uint32_t)(c>>16&0xff)+1) >> 1 <<16 |
    //    ((uint32_t)(b>>24&0xff)+(uint32_t)(c>>24&0xff)+1) >> 1 <<24;
}
使用道具 举报 回复 支持 反对
发表于 2013-11-21 12:38:41
afterglow 发表于 2013-11-21 12:37
我确实没有说谎,为了确定你说的结论,我再一次编译(没写-arch 或者补上 -arch compute_10)确实生成了pt ...

哦。版主说错了。

这种情况是正常的,因为ptx不是生成的,是你嵌入的。

此时将不检查任何错误(包括语法错误)(甚至你可以潜入一条afterglow.s32指令)

深表歉意。
使用道具 举报 回复 支持 反对
12下一页
发新帖
您需要登录后才可以回帖 登录 | 立即注册