LLVM NVPTX 백엔드 struct parameter zero size

Question

LLVM의 NVPTX 백엔드로 생성 된 PTX 어셈블리를로드 할 때 확실하지 않은 예외가 발생합니다. (나는 ManagedCuda에서 PTX로드 해요 - http://managedcuda.codeplex.com/를)

여기

ErrorNoBinaryForGPU: This indicates that there is no kernel image available that is suitable for the device. This can occur when a user specifies code generation options for a particular CUDA source file that do not include the corresponding device configuration. 

모듈에 대한 LLVM IR입니다

; ModuleID = 'Module' 
target triple = "nvptx64-nvidia-cuda" 

%testStruct = type { i32 } 

define void @kernel(i32 addrspace(1)*) { 
entry: 
    %1 = alloca %testStruct 
    store %testStruct zeroinitializer, %testStruct* %1 
    %2 = load %testStruct* %1 
    call void @structtest(%testStruct %2) 
    ret void 
} 

define void @structtest(%testStruct) { 
entry: 
    ret void 
} 

!nvvm.annotations = !{!0} 

!0 = metadata !{void (i32 addrspace(1)*)* @kernel, metadata !"kernel", i32 1} 

을 (그것은이 도구에 의해 생성 된 이후 약간 이상해) 여기에있다 PTX

// 
// Generated by LLVM NVPTX Back-End 
// 

.version 3.1 
.target sm_20 
.address_size 64 

     // .globl  kernel 
.visible .func structtest 
(
     .param .b0 structtest_param_0 
) 
; 

.visible .entry kernel(
     .param .u64 kernel_param_0 
) 
{ 
     .local .align 8 .b8  __local_depot0[8]; 
     .reg .b64  %SP; 
     .reg .b64  %SPL; 
     .reg .s32  %r<2>; 
     .reg .s64  %rl<2>; 

     mov.u64   %rl1, __local_depot0; 
     cvta.local.u64 %SP, %rl1; 
     mov.u32   %r1, 0; 
     st.u32 [%SP+0], %r1; 
     // Callseq Start 0 
     { 
     .reg .b32 temp_param_reg; 
     // <end>} 
     .param .align 4 .b8 param0[4]; 
     st.param.b32 [param0+0], %r1; 
     call.uni 
     structtest, 
     (
     param0 
     ); 

     //{ 
     }// Callseq End 0 
     ret; 
} 

     // .globl  structtest 
.visible .func structtest(
     .param .b0 structtest_param_0 
) 
{ 


     ret; 
} 

를 결과 어떻게 PTX를 읽는 아무 생각이 없다,하지만 난 문제를 느끼는 것은의 .b0 비트과 관련이있다 structtest 함수 정의에서 0.

구조체가 아닌 값 (정수 또는 포인터 등)을 전달하면 올바르게 작동하고 .b0이 작동합니다. 비트는 .b32 또는 .b64과 같은 정상적인 것을 읽습니다. 트리플 nvptx - 엔비디아 CUDA를 위해 (32 비트) 수행 아무 것도 변경하지뿐만 아니라 데이터 레이아웃 제외/포함

가

http://llvm.org/docs/NVPTXUsage.html

에서 제안이는 NVPTX 백엔드 버그, 아니면 내가 뭔가 잘못하고있는 중이 야 ?

---

업데이트 :

나는이를 찾고 있어요 - http://llvm.org/docs/doxygen/html/NVPTXAsmPrinter_8cpp_source.html -와 유형이 01568 라인을 통해 떨어지는 경우, 분명히 원시적 형이 아니며, Ty->getPrimitiveSizeInBits()는 0을 반환로 나타납니다. (어쨌든 적어도 내 추측이다)

구조체인지 확인하고 주소를 가져 와서 인수를 byval으로 만들고 이후에 구조체를 역 참조하기 위해 특수 케이스를 추가해야합니까? 해키 한 솔루션처럼 보이지만 해결 방법을 모르겠습니다.

Answer 1

컴파일에서 오류 메시지 버퍼를 가져 오려고 했습니까? 내가 말하는 당신의 PTX를 실행하면

CudaContext ctx = new CudaContext(); 
CudaJitOptionCollection options = new CudaJitOptionCollection(); 
CudaJOErrorLogBuffer err = new CudaJOErrorLogBuffer(1024); 
options.Add(err); 
try 
{ 
    ctx.LoadModulePTX("test.ptx", options); 
} 
catch 
{ 
    options.UpdateValues(); 
    MessageBox.Show(err.Value);     
} 

: managedCuda에서이 같은 일 것입니다

ptxas 응용 프로그램 PTX 입력, 라인 (12); 치명적인 '.b0'근처에 구문 분석 오류 : 구문 오류

ptxas 치명적인 :.! 너무 큰 어셈블리 인해 B0와 추측을 지원하는 오류 "

에 중단