修改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上。
Comments