PTXPlus笔记

修改ptxplus建议

在gpgpusim源码src/cuda-sim/ptx_loader.cc的函数

char *gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptxfilename, const std::string elffilename,const std::string sassfilename)中,std::ifstream fileStream(fname_ptxplus, std::ios::in);一行之前添加内容如下:

    // FindHao 在这里添加对ptxplus文件的处理
    char commandline2[1024];
    snprintf(commandline2, 1024, "python3 ptxplus_plus.py %s",fname_ptxplus);
    fflush(stdout);
    result = system(commandline2);
    if (result) {
        printf("FindHao: ERROR ** could not execute %s\n", commandline2);
        exit(1);
    }


	// 在这行之前
    // Get ptxplus from file
    std::ifstream fileStream(fname_ptxplus, std::ios::in);
    std::string text, line;

再新建ptxplus_plus.py来修改ptxplus文件即可。

如果代码和编译参数不变,ptxplus文件是不变的,python脚本中可以直接用改后的ptxplus文件覆盖新的ptxplus文件即可。

ptxplus语法特殊的地方

Kernel源码:

__global__ void test(int* aa, float* bb){
    int thx = threadIdx.x;
    aa[thx] = thx + 1;
}

对应的的ptxplus源码

.entry  _Z4testPiPf (
	.param  .u64 __cudaparm__Z4testPiPf_aa ,
	.param  .u64 __cudaparm__Z4testPiPf_bb )
{
	.reg .u32 $r<3>;

	
	cvt.u32.u16 $r0, $r0.lo;
	shl.u32 $r1, $r0, 0x00000002;
	add.u32 $r2, $r0, 0x00000001;
	add.u32 $r0, s[0x0010], $r1;
	st.global.u32 [$r0], $r2;

	l_exit: exit;
}

可以看到,默认r0.lo存储的是threadid.x,传入参数的起始地址则存储在了shared memory上。

文章版权归 FindHao 所有丨本站默认采用CC-BY-NC-SA 4.0协议进行授权|
转载必须包含本声明,并以超链接形式注明作者 FindHao 和本文原始地址:
https://www.findhao.net/easycoding/2477

你可能喜欢:

Find

新浪微博(FindHaoX86)QQ群:不安分的Coder(375670127)不安分的Coder 微信公众号(findhao-net)

发表评论

电子邮件地址不会被公开。 必填项已用*标注