我試圖找到一種方法，從CUDA C.調用函數PTX（.FUNC） 說我有一個PTX功能是這樣的：如何從CUDA C調用ptx函數？
.func (.reg .s32 %res) inc_ptr (.reg .s32 %ptr, .reg .s32 %inc)
{
add.s32 %res, %ptr, %inc;
ret;
}
我知道我可以從PTX調用它像這樣：
call (%d), inc_ptr, (%s, %d);
但我不知道如何將它從CUDA C. 打電話，我知道我可以內聯彙編PTX與ASM（），但我還沒有找到一種方法，內聯函數。 希望有人能幫助！
謝謝！
據我所知，CUDA C支持asm，安裝cuda工具包後有一個doc文件位於doc目錄下。
是CUDA C支持asm，而且正如我寫的，我知道我可以在CUDA C中內聯彙編代碼，但我不知道如何內聯ptx彙編函數（.func）。 – fursund
這可以使用CUDA 5.0中引入的單獨編譯工具來完成。我不相信在「整個」程序編譯模式或CUDA 5.0之前的工具包版本或3.1之前的PTX修訂版本中，都有這樣的方法。
它可能是最容易說明如何用一個工作的例子做到這一點。讓我們先從一個簡單的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運行的設備代碼鏈接階段（至少在CUDA 5中）期望所有設備代碼都存在於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
/*0008*/ /*0x0040dc8580000000*/ LD R3, [R4];
/*0010*/ /*0x00301c8580000000*/ LD R0, [R3];
/*0018*/ /*0x14001c0348000000*/ IADD R0, R0, R5;
/*0020*/ /*0x00301c8590000000*/ ST [R3], R0;
/*0028*/ /*0x00001de790000000*/ RET;
/*0030*/ /*0x00001de440000000*/ NOP CC.T;
/*0038*/ /*0x00001de440000000*/ NOP CC.T;
/*0040*/ /*0xe0001de74003ffff*/ BRA 0x40;
/*0048*/ /*0x00001de440000000*/ NOP CC.T;
/*0050*/ /*0x00001de440000000*/ NOP CC.T;
/*0058*/ /*0x00001de440000000*/ NOP CC.T;
/*0060*/ /*0x00001de440000000*/ NOP CC.T;
/*0068*/ /*0x00001de440000000*/ NOP CC.T;
/*0070*/ /*0x00001de440000000*/ NOP CC.T;
/*0078*/ /*0x00001de440000000*/ 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
/*0008*/ /*0x10005de428004001*/ MOV R1, c [0x0] [0x44];
/*0010*/ /*0x20105d034800c000*/ IADD R1, R1, -0x8;
/*0018*/ /*0x00019de428004005*/ MOV R6, c [0x0] [0x140];
/*0020*/ /*0x10101c034800c000*/ IADD R0, R1, 0x4;
/*0028*/ /*0x8400dc042c000000*/ S2R R3, SR_Tid_X;
/*0030*/ /*0x90041c0348004000*/ IADD R16, R0, c [0x0] [0x24];
/*0038*/ /*0x94001c042c000000*/ S2R R0, SR_CTAid_X;
/*0048*/ /*0xd0009de428004000*/ MOV R2, c [0x0] [0x34];
/*0050*/ /*0x91045d0348004000*/ IADD R17, R16, -c [0x0] [0x24];
/*0058*/ /*0x40011de428000000*/ MOV R4, R16;
/*0060*/ /*0xa0015ca320064000*/ IMAD R5, R0, c [0x0] [0x28], R3;
/*0068*/ /*0x01119c85c8000000*/ STL [R17], R6;
/*0070*/ /*0xa0209ca350004000*/ IMUL R2, R2, c [0x0] [0x28];
/*0078*/ /*0x0001000710000000*/ JCAL 0x0;
/*0088*/ /*0x0110dc85c0000000*/ LDL R3, [R17];
/*0090*/ /*0x20001de428004005*/ MOV R0, c [0x0] [0x148];
/*0098*/ /*0x00049c4340004005*/ ISCADD R18, R0, c [0x0] [0x140], 0x2;
/*00a0*/ /*0x4831dc031b0e0000*/ ISETP.GE.U32.AND P0, pt, R3, R18, pt;
/*00a8*/ /*0x000001e780000000*/ @P0 EXIT;
/*00b0*/ /*0x1004dde428004005*/ MOV R19, c [0x0] [0x144];
/*00b8*/ /*0x0034dc8590000000*/ ST [R3], R19;
/*00c8*/ /*0x40011de428000000*/ MOV R4, R16;
/*00d0*/ /*0x08015de428000000*/ MOV R5, R2;
/*00d8*/ /*0x0001000710000000*/ JCAL 0x0;
/*00e0*/ /*0x0110dc85c0000000*/ LDL R3, [R17];
/*00e8*/ /*0x4831dc03188e0000*/ ISETP.LT.U32.AND P0, pt, R3, R18, pt;
/*00f0*/ /*0x000001e74003ffff*/ @P0 BRA 0xb8;
/*00f8*/ /*0x00001de780000000*/ EXIT;
/*0100*/ /*0xe0001de74003ffff*/ BRA 0x100;
/*0108*/ /*0x00001de440000000*/ NOP CC.T;
/*0110*/ /*0x00001de440000000*/ NOP CC.T;
/*0118*/ /*0x00001de440000000*/ NOP CC.T;
/*0120*/ /*0x00001de440000000*/ NOP CC.T;
/*0128*/ /*0x00001de440000000*/ NOP CC.T;
/*0130*/ /*0x00001de440000000*/ NOP CC.T;
/*0138*/ /*0x00001de440000000*/ NOP CC.T;
.....................................
Function : inc_ptr
/*0008*/ /*0x0040dc8580000000*/ LD R3, [R4];
/*0010*/ /*0x00301c8580000000*/ LD R0, [R3];
/*0018*/ /*0x14001c0348000000*/ IADD R0, R0, R5;
/*0020*/ /*0x00301c8590000000*/ ST [R3], R0;
/*0028*/ /*0x00001de790000000*/ RET;
/*0030*/ /*0x00001de440000000*/ NOP CC.T;
/*0038*/ /*0x00001de440000000*/ NOP CC.T;
/*0040*/ /*0xe0001de74003ffff*/ BRA 0x40;
/*0048*/ /*0x00001de440000000*/ NOP CC.T;
/*0050*/ /*0x00001de440000000*/ NOP CC.T;
/*0058*/ /*0x00001de440000000*/ NOP CC.T;
/*0060*/ /*0x00001de440000000*/ NOP CC.T;
/*0068*/ /*0x00001de440000000*/ NOP CC.T;
/*0070*/ /*0x00001de440000000*/ NOP CC.T;
/*0078*/ /*0x00001de440000000*/ NOP CC.T;
........................
有可能是可與工具鏈播放達到這個其他的招數，但這種做法肯定工作。
增加了一個答案，如果有人願意讚揚它並/或接受它，那麼可以從未回答的問題列表中獲得此答案。 – talonmies
做得很好，天才！ –
好的一個問題。如何創建一個包含設備鏈接對象（memset.out）的可執行文件？ –
不幸的是，我認爲這可能是不可能的。問題在於CUDA C沒有提供設備代碼的鏈接器。因此，從一個內核運行的所有內容都必須位於同一個.cu文件中。我希望我錯了，因爲IMO，NVIDIA在內聯PTX的語法方面做得不好。 –
@RogerDahl是的，這也是我所害怕的。但是看起來CUDA 5中可能會出現設備代碼鏈接器（http://developer.download.nvidia.com/assets/cuda/files/CUDADownloads/GPU_Library_Object_Linking.pdf）。儘管如此，我還沒有發現如何在CUDA 5 Preview中做到這一點。 – fursund