我想用nvrtc JIT编译器来编译CUDA内核来提高我的应用程序的性能(所以我有更多的指令提取但我保存了多个数组访问)。cuModuleGetFunction返回未找到
这些函数看起来像例如这样,由我的函数发生器(不那么重要)产生:
extern "C" __device__ void GetSumOfBranches(double* branches, double* outSum)
{
double sum = (branches[38])+(-branches[334])+(-branches[398])+(-branches[411]);
*outSum = sum;
}
我编译上面用下面的函数代码:
CUfunction* FunctionGenerator::CreateFunction(const char* programText)
{
// When I comment this statement out the output of the PTX file is changing
// what is the reson?!
// Bug?
std::string savedString = std::string(programText);
nvrtcProgram prog;
nvrtcCreateProgram(&prog, programText, "GetSumOfBranches.cu", 0, NULL, NULL);
const char *opts[] = {"--gpu-architecture=compute_52", "--fmad=false"};
nvrtcCompileProgram(prog, 2, opts);
// Obtain compilation log from the program.
size_t logSize;
nvrtcGetProgramLogSize(prog, &logSize);
char *log = new char[logSize];
nvrtcGetProgramLog(prog, log);
// Obtain PTX from the program.
size_t ptxSize;
nvrtcGetPTXSize(prog, &ptxSize);
char *ptx = new char[ptxSize];
nvrtcGetPTX(prog, ptx);
printf("%s", ptx);
CUdevice cuDevice;
CUcontext context;
CUmodule module;
CUfunction* kernel;
kernel = (CUfunction*)malloc(sizeof(CUfunction));
cuInit(0);
cuDeviceGet(&cuDevice, 0);
cuCtxCreate(&context, 0, cuDevice);
auto resultLoad = cuModuleLoadDataEx(&module, ptx, 0, 0, 0);
auto resultGetF = cuModuleGetFunction(kernel, module, "GetSumOfBranches");
return kernel;
}
一切工作,除了cuModuleGetFunction
正在恢复CUDA_ERROR_NOT_FOUND
。出现此错误的原因是在PTX文件中找不到GetSumOfBranches
。
然而printf("%s", ptx);
输出是这样的:
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-19856038
// Cuda compilation tools, release 7.5, V7.5.17
// Based on LLVM 3.4svn
//
.version 4.3
.target sm_52
.address_size 64
// .globl GetSumOfBranches
.visible .func GetSumOfBranches(
.param .b64 GetSumOfBranches_param_0,
.param .b64 GetSumOfBranches_param_1
)
{
.reg .f64 %fd<8>;
.reg .b64 %rd<3>;
ld.param.u64 %rd1, [GetSumOfBranches_param_0];
ld.param.u64 %rd2, [GetSumOfBranches_param_1];
ld.f64 %fd1, [%rd1+304];
ld.f64 %fd2, [%rd1+2672];
sub.rn.f64 %fd3, %fd1, %fd2;
ld.f64 %fd4, [%rd1+3184];
sub.rn.f64 %fd5, %fd3, %fd4;
ld.f64 %fd6, [%rd1+3288];
sub.rn.f64 %fd7, %fd5, %fd6;
st.f64 [%rd2], %fd7;
ret;
}
在我optinion一切都很好,并GetSumOfBranches
通过前人的精力来cuModuleGetFunction
发现。你能解释我为什么吗?
第二个问题
当我outcomment std::string savedString = std::string(programText);
那么PTX的输出就是:
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-19856038
// Cuda compilation tools, release 7.5, V7.5.17
// Based on LLVM 3.4svn
//
.version 4.3
.target sm_52
.address_size 64
因为savedString
完全不使用这是奇怪...
为什么你有标记为'__device__'的问题?这不是CUDA内核。你不能从主机代码中调用它。无论如何,对于类型问题“为什么这不起作用?”因此你应该提供一个[mcve] –
我没有想过我只编译内核..我认为我也可以编译设备函数。你刚刚解决了我的问题(也许你想回答我的问题以获得积分? - 否则谢谢你)。在下一个问题中,我将提供一个最小的,完整的和可验证的例子。 – Jens