CUDA C에서 ptx 함수를 호출하려면 어떻게해야합니까?

Question

내가 내가 이런 PTX 기능을 가지고 말 CUDA C.에서 PTX 기능 (.func)를 호출 할 수있는 방법을 찾기 위해 노력하고 있어요 : 나는과 같이 PTX에서 호출 할 수 있습니다 알고

.func (.reg .s32 %res) inc_ptr (.reg .s32 %ptr, .reg .s32 %inc) 
{ 
    add.s32 %res, %ptr, %inc; 
    ret; 
} 

를 :

call (%d), inc_ptr, (%s, %d); 

는하지만 내가 ASM()와 PTX 어셈블리를 인라인 수있는 방법을 알고 CUDA C. 에서 호출하는 아무 생각이 없다,하지만 난 함수를 인라인 할 수있는 방법을 발견하지 않았습니다. 누군가 도움을 줄 수 있기를 바랍니다.

감사합니다.

Answer 1

내가 아는 한, CUDA C는 asm을 지원합니다. cuda tool kit를 설치 한 후 doc 디렉토리에있는 doc이 있습니다.

Answer 2

이것은 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; 
     ........................ 

조립 된 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; 
     ........................ 

이를 달성하기 위해 툴체인과 함께 연주 할 수있는 다른 트릭이있을 수도 있지만,이 방법은 확실히 작동합니다.