可以使用CUDA 5.0引入的单独编译功能来实现这一点。我认为在“整个”程序编译模式或CUDA 5.0以前的工具包版本或PTX修订版3.1之前,无法做到这一点。
用一个简单的PTX函数为例,说明如何实现这一点,它类似于您的示例:增加指针。
.version 3.1
.target sm_30
.address_size 32
.visible .func inc_ptr(.param .b32 ptr, .param .b32 inc)
{
.reg .s32 %r<6>
ld.param.u32 %r1, [ptr]
ld.param.u32 %r2, [inc]
ld.u32 %r3, [%r1]
ld.u32 %r4, [%r3]
add.s32 %r5, %r4, %r2
st.u32 [%r3], %r5
ret
}
可以使用ptxas
将此内容编译为可重定位设备对象,然后打包到fatbinary容器文件中。后面这个步骤似乎是至关重要的。默认的ptxas
输出只是一个可重定位的elf
对象,没有生成任何fatbinary容器。似乎nvcc运行的设备代码链接阶段(至少在CUDA5中)期望所有设备代码都在fatbinary容器中存在。否则,链接将失败。结果看起来像这样:
$ ptxas -arch=sm_30 -c -o inc_ptr.gpu.o inc_ptr.ptx
$ fatbinary -arch=sm_30 -create inc_ptr.fatbin -elf inc_ptr.gpu.o
$ cuobjdump -sass inc_ptr.fatbin
Fatbin elf code:
================
arch = sm_30
code version = [1,6]
producer = <unknown>
host = mac
compile_size = 32bit
code for sm_30
Function : inc_ptr
LD R3, [R4];
LD R0, [R3];
IADD R0, R0, R5;
ST [R3], R0;
RET;
NOP CC.T;
NOP CC.T;
BRA 0x40;
NOP CC.T;
NOP CC.T;
NOP CC.T;
NOP CC.T;
NOP CC.T;
NOP CC.T;
NOP CC.T;
........................
您可以看到fatbinary包含了从组装的PTX中提取出来的微码。有了准备好的设备函数fatbin,您可以在CUDA C代码中执行以下操作:
extern "C" __device__ void inc_ptr(int* &ptr, const int inc);
__global__
void memsetkernel(int *inout, const int val, const int N)
{
int stride = blockDim.x * gridDim.x;
int *p = inout;
inc_ptr(p, threadIdx.x + blockDim.x*blockIdx.x);
for(; p < inout+N; inc_ptr(p, stride)) *p = val;
}
int main(void)
{
const int n=10;
int *p;
cudaMalloc((void**)&p, sizeof(int)*size_t(n));
memsetkernel<<<1,32>>>(p, 5, n);
return 0;
}
在单独编译模式下,设备代码工具链将遵循
extern
声明,并且(只要您控制好符号重整),设备函数 fatbinary 可以与其他设备和主机代码链接以产生最终对象。
$ nvcc -arch=sm_30 -Xptxas="-v" -dlink -o memset.out inc_ptr.fatbin memset_kernel.cu
ptxas info : 0 bytes gmem
ptxas info : Compiling entry function '_Z12memsetkernelPiii' for 'sm_30'
ptxas info : Function properties for _Z12memsetkernelPiii
8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 20 registers, 332 bytes cmem[0]
$ cuobjdump -sass memset.out
Fatbin elf code:
================
arch = sm_30
code version = [1,6]
producer = <unknown>
host = mac
compile_size = 32bit
identifier = inc_ptr.fatbin memset_kernel.cu
code for sm_30
Function : _Z12memsetkernelPiii
MOV R1, c [0x0] [0x44];
IADD R1, R1, -0x8;
MOV R6, c [0x0] [0x140];
IADD R0, R1, 0x4;
S2R R3, SR_Tid_X;
IADD R16, R0, c [0x0] [0x24];
S2R R0, SR_CTAid_X;
MOV R2, c [0x0] [0x34];
IADD R17, R16, -c [0x0] [0x24];
MOV R4, R16;
IMAD R5, R0, c [0x0] [0x28], R3;
STL [R17], R6;
IMUL R2, R2, c [0x0] [0x28];
JCAL 0x0;
LDL R3, [R17];
MOV R0, c [0x0] [0x148];
ISCADD R18, R0, c [0x0] [0x140], 0x2;
ISETP.GE.U32.AND P0, pt, R3, R18, pt;
@P0 EXIT;
MOV R19, c [0x0] [0x144];
ST [R3], R19;
MOV R4, R16;
MOV R5, R2;
JCAL 0x0;
LDL R3, [R17];
ISETP.LT.U32.AND P0, pt, R3, R18, pt;
@P0 BRA 0xb8;
EXIT;
BRA 0x100;
NOP CC.T;
NOP CC.T;
NOP CC.T;
NOP CC.T;
NOP CC.T;
NOP CC.T;
NOP CC.T;
.....................................
Function : inc_ptr
LD R3, [R4];
LD R0, [R3];
IADD R0, R0, R5;
ST [R3], R0;
RET;
NOP CC.T;
NOP CC.T;
BRA 0x40;
NOP CC.T;
NOP CC.T;
NOP CC.T;
NOP CC.T;
NOP CC.T;
NOP CC.T;
NOP CC.T;
........................
可能还有其他技巧可以使用工具链来实现这一点,但这种方法肯定有效。