문제

2014-07-21 4 views
8

다음과 같이 나는, 생성자의 커널을 호출하는 클래스가 :문제

"ScalarField.h"

#include <iostream> 

    void ERROR_CHECK(cudaError_t err,const char * msg) { 
     if(err!=cudaSuccess) { 
      std::cout << msg << " : " << cudaGetErrorString(err) << std::endl; 
      std::exit(-1); 
     } 
    } 

    class ScalarField { 
    public: 
     float* array; 
     int dimension; 

     ScalarField(int dim): dimension(dim) { 
      std::cout << "Scalar Field" << std::endl; 
      ERROR_CHECK(cudaMalloc(&array, dim*sizeof(float)),"cudaMalloc"); 
     } 
    }; 

"classA.h"

#include "ScalarField.h" 


static __global__ void KernelSetScalarField(ScalarField v) { 
    int index = threadIdx.x + blockIdx.x * blockDim.x; 
    if (index < v.dimension) v.array[index] = 0.0f; 
} 

class A { 
public: 
    ScalarField v; 

    A(): v(ScalarField(3)) { 
     std::cout << "Class A" << std::endl; 
     KernelSetScalarField<<<1, 32>>>(v); 
     ERROR_CHECK(cudaGetLastError(),"Kernel"); 
    } 
}; 

"main.cu"

#include "classA.h" 

A a_object; 

int main() { 
    std::cout << "Main" << std::endl; 
    return 0; 
} 

main (A a_object;)에서이 클래스를 인스턴스화하면 오류가 발생하지 않습니다. 그러나 메인 (main) 외부에서 인스턴스를 생성하면 커널 정의시 (class A {...} a_object;) 커널이 시작될 때 "invalid device function"오류가 발생합니다. 왜 그런 일이 일어날까요?

편집

업데이트 코드는보다 완전한 예제를 제공합니다. Raxvan에 의해 주석의 조언에 따라 편집 2

, 나는 내가 주요 외부 ScalarField 생성자에서 사용되는 dimensions 변수도 (다른 클래스에서) 정의를 말하고 싶었지만, 다른 모든 전에. 그것은 설명 일 수 있습니까? 디버거는 dimensions에 대한 올바른 값을 표시하고있었습니다.

+0

다음 질문에 대한 답변을 제공 할 수 있습니다 : 클래스 A는 자체 파일이지만 커널은 파일 확장자입니다. 다른 사람이 문제를 복제 할 수 있도록 충분한 코드를 제공해야합니다. – deathly809

+4

@Noel Perez Gonzalez'a_Object'를 전역 변수로 정의하면 전역 데이터 초기화 중에 실행을 시작합니다. 이것은 실행 순서를 알 수있는 방법이 없으므로 매우 나쁜 관행입니다. 이를 염두에두면 모든 CUDA를 초기화하는 코드가 글로벌 데이터보다 나중에 실행될 수 있습니다. – Raxvan

+0

더 많은 코드로 질문을 업데이트했습니다 (컴파일하지 않았 음에 유의하십시오). @Raxvan 조언을 주셔서 감사합니다, 나는 방금 런타임 순서와 컴파일 순서가 같다고 생각했습니다. – Noel

답변

12

짧은 버전 :

class A 메인 외부 인스턴스화하는 문제에 대한 근본적인 이유는 커널로 CUDA 런타임 라이브러리를 초기화하는 데 필요한 특정 후크 루틴이 먼저 실행되지 않을 것입니다 class A의 생성자가 호출되고 있습니다. 이것은 정적 객체가 C++ 실행 모델에서 인스턴스화되고 초기화되는 순서에 대한 보장이 없기 때문에 발생합니다. 전역 범위 클래스는 CUDA 설정을 수행하는 전역 범위 개체가 초기화되기 전에 인스턴스화됩니다. 커널 코드는 호출되기 전에 컨텍스트에로드되지 않으며 런타임 오류가 발생합니다.

내가 알 수있는 한, 이것은 CUDA 런타임 API의 진정한 한계이며 사용자 코드에 쉽게 고정되어있는 것이 아닙니다. 간단한 예제에서는 커널 호출을 cudaMemset 또는 비 심볼 기반 런타임 API memset 함수 중 하나를 호출하여 대체 할 수 있습니다. 이 문제는 런타임에 런타임 API를 통해로드 된 사용자 커널이나 장치 기호로 완전히 제한됩니다. 이런 이유로, 빈 디폴트 생성자는 또한 당신의 문제를 해결할 것이다. 디자인의 관점에서 보았을 때, 나는 생성자에서 커널을 호출하는 패턴에 대해 매우 모호 할 것이다. 기본 생성자 또는 소멸자에 의존하지 않는 GPU 설정/분해 클래스에 대한 특정 메소드를 추가하는 것이 훨씬 깔끔하고 오류가 발생하기 쉬운 디자인 인 IMHO가됩니다. 상세

: 어떤 런타임 API 프로그램의 fatbin 페이로드에 포함 된 커널, 질감과 정적으로 정의 장치 기호를로드 및 등록을 실행해야합니다 내부적으로 생성 루틴 (__cudaRegisterFatBinary가) 있습니다

커널 이전에 CUDA 드라이버 API를 오류없이 호출 할 수 있습니다. 이것은 런타임 API의 "게으른"컨텍스트 초기화 기능의 일부입니다.다음과 같이 직접 확인할 수 있습니다.

게시 한 수정 된 예제의 gdb 추적입니다. 인 (여기

[email protected]:~$ gdb a.out 
GNU gdb (Ubuntu/Linaro 7.4-2012.04-0ubuntu2.1) 7.4-2012.04 
Copyright (C) 2012 Free Software Foundation, Inc. 
License GPLv3+: GNU GPL version 3 or later <http://gnu.org/licenses/gpl.html> 
This is free software: you are free to change and redistribute it. 
There is NO WARRANTY, to the extent permitted by law. Type "show copying" 
and "show warranty" for details. 
This GDB was configured as "x86_64-linux-gnu". 
For bug reporting instructions, please see: 
<http://bugs.launchpad.net/gdb-linaro/>... 
Reading symbols from /home/talonmies/a.out...done. 
(gdb) break '__cudaRegisterFatBinary' 
Breakpoint 1 at 0x403180 
(gdb) run 
Starting program: /home/talonmies/a.out 
[Thread debugging using libthread_db enabled] 
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1". 
Scalar Field 
[New Thread 0x7ffff5a63700 (LWP 10774)] 
Class A 
Kernel : invalid device function 
[Thread 0x7ffff5a63700 (LWP 10774) exited] 
[Inferior 1 (process 10771) exited with code 0377] 

이 같은 절차 main 내부 A 인스턴스화이 시간 : 나는 __cudaRegisterFatBinary에 중단 점을 삽입하고 정적 A 생성자를 호출하기 전에 도달되지 않고 커널 발사가 실패합니다 게으른 설치를 수행하는 객체가 초기화 된 후 보장)가 발생합니다 :

[email protected]:~$ cat main.cu 
#include "classA.h" 


int main() { 
    A a_object; 
    std::cout << "Main" << std::endl; 
    return 0; 
} 

[email protected]:~$ nvcc --keep -arch=sm_30 -g main.cu 
[email protected]:~$ gdb a.out 
GNU gdb (Ubuntu/Linaro 7.4-2012.04-0ubuntu2.1) 7.4-2012.04 
Copyright (C) 2012 Free Software Foundation, Inc. 
License GPLv3+: GNU GPL version 3 or later <http://gnu.org/licenses/gpl.html> 
This is free software: you are free to change and redistribute it. 
There is NO WARRANTY, to the extent permitted by law. Type "show copying" 
and "show warranty" for details. 
This GDB was configured as "x86_64-linux-gnu". 
For bug reporting instructions, please see: 
<http://bugs.launchpad.net/gdb-linaro/>... 
Reading symbols from /home/talonmies/a.out...done. 
(gdb) break '__cudaRegisterFatBinary' 
Breakpoint 1 at 0x403180 
(gdb) run 
Starting program: /home/talonmies/a.out 
[Thread debugging using libthread_db enabled] 
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1". 

Breakpoint 1, 0x0000000000403180 in __cudaRegisterFatBinary() 
(gdb) cont 
Continuing. 
Scalar Field 
[New Thread 0x7ffff5a63700 (LWP 11084)] 
Class A 
Main 
[Thread 0x7ffff5a63700 (LWP 11084) exited] 
[Inferior 1 (process 11081) exited normally] 

을이 정말 당신을위한 타격 문제가 있다면, 엔비디아의 개발자 지원 센터에 문의하고, 버그 리포트를 제기 건의 할 것입니다.

+0

우수 답변. "세계적으로"초기화 된 Thrust 객체에 대해서도 같은 결과가 발생합니까? – JackOLantern

+0

매우 유익한 답변입니다. 내가 추천 한대로 데이터를 초기화하는 멤버 함수에 의지했다. 감사. – Noel