__CUDA_ARCH__ und Kernel-Aufruf in __host__ __device__ Funktion

Question

ich auf eine seltsame Wirkung kam:

#define CUDA_ERR_CHECK(call) call 

#include <assert.h> 
#include <iostream> 

using namespace std; 

#if defined(__CUDACC__) 

// Determine the size of type on device. 
template<typename T> 
__global__ void deviceSizeOf(size_t* result) 
{ 
    *result = sizeof(T); 
} 

// Device memory aligned vector. 
template<typename T> 
class VectorDevice 
{ 
    T* data; 
    size_t size; 
    int dim, dim_aligned; 

public : 
    __host__ __device__ 
    VectorDevice() : data(NULL), size(0), dim(0) { } 

    __host__ __device__ 
    VectorDevice(int dim_) : data(NULL), size(0), dim(dim_) 
    { 
     dim_aligned = dim_; 
     if (dim_ % AVX_VECTOR_SIZE) 
      dim_aligned = dim + AVX_VECTOR_SIZE - dim_ % AVX_VECTOR_SIZE; 
#if !defined(__CUDA_ARCH__) 
     // Determine the size of target type. 
     size_t size, *dSize; 
     CUDA_ERR_CHECK(cudaMalloc(&dSize, sizeof(size_t))); 
     deviceSizeOf<T><<<1, 1>>>(dSize); 
     CUDA_ERR_CHECK(cudaGetLastError()); 
     CUDA_ERR_CHECK(cudaDeviceSynchronize()); 
     CUDA_ERR_CHECK(cudaMemcpy(&size, dSize, sizeof(size_t), cudaMemcpyDeviceToHost)); 
     CUDA_ERR_CHECK(cudaFree(dSize)); 

     // Make sure the size of type is the same on host and on device. 
     if (size != sizeof(T)) 
     { 
      std::cerr << "Unexpected unequal sizes of type T in VectorDevice<T> on host and device" << std::endl; 
      exit(2); 
     } 
#endif 
    } 
}; 

#endif // __CUDACC__ 

int main() 
{ 
    VectorDevice<int> v(10); 

    return 0; 
} 

Hier wird ein Kernel wird von der Host-Version von __host__ __device__ Konstruktor aufgerufen. Überraschenderweise wenn Sie diesen Code ausgeführt verlässt er leise mit Code 1 vom Kernel Aufruf Wrapper:

(gdb) make 
nvcc -arch=sm_30 test.cu -o test -DAVX_VECTOR_SIZE=32 
(gdb) b exit 
Breakpoint 1 at 0x7ffff711b1e0: file exit.c, line 104. 
(gdb) r 
Breakpoint 1, __GI_exit (status=1) at exit.c:104 
104 exit.c: No such file or directory. 
(gdb) f 3 
#3 0x0000000000402c36 in VectorDevice<int>::VectorDevice(int)() 
(gdb) f 2 
#2 0x0000000000402cb0 in void deviceSizeOf<int>(unsigned long*)() 
(gdb) f 1 
#1 0x0000000000402ad2 in void __wrapper__device_stub_deviceSizeOf<int>(unsigned long*&)() 
(gdb) disass 
Dump of assembler code for function _Z35__wrapper__device_stub_deviceSizeOfIiEvRPm: 
    0x0000000000402abc <+0>: push %rbp 
    0x0000000000402abd <+1>: mov %rsp,%rbp 
    0x0000000000402ac0 <+4>: sub $0x10,%rsp 
    0x0000000000402ac4 <+8>: mov %rdi,-0x8(%rbp) 
    0x0000000000402ac8 <+12>: mov $0x1,%edi 
    0x0000000000402acd <+17>: callq 0x402270 <exit@plt> 
End of assembler dump. 

Weitere Untersuchungen geht hervor, dass der Kernel-Code nicht in Cubin angezeigt wird, und dass __CUDA_ARCH__ irgendwie in dieses Verhalten beteiligt.

So, 2 Fragen:

1) Warum passiert das das?

2) Wie wird __CUDA_ARCH__ für bedingte Kompilierung von __host__ __device__ Code in Kombination mit Host-Side-Kernel-Aufrufe verwendet?

Danke!

UPDATE: Das gleiche Beispiel wird in Abschnitt E.2.2.1 Punkt 2 des C-Programmierhandbuchs gezeigt. Es ist jedoch immer noch unklar, wie man dieses Problem richtig behandelt.

Answer 1

1) Warum passiert das?

Es geschieht, weil Sie auf the specific restriction verstärken Sie in der Programmieranleitung hingewiesen: die Templat-Instanziierung deviceSizeOf für <int> muss sowohl bei __CUDA_ARCH__ definiert ist und wenn es nicht definiert ist. Wenn Sie eingeschränkte Formulare verwenden, ist das Verhalten nicht definiert.

2) Wie __CUDA_ARCH__ für die bedingte Kompilierung von __host__ __device__ Code in Kombination mit hostseitigen Kernel Anrufen zu benutzen?

Ein möglicher Ansatz wäre die Instanziierung der Kernfunktion für den Typen zu erzwingen <int> unabhängig von dem __CUDA_ARCH__ Makro.

Sie können dies tun, indem Sie die folgende Zeile unmittelbar nach der Kernel-Template-Definition hinzu:

template __global__ void deviceSizeOf<int>(size_t *); 

Wenn ich diese Zeile nach dem Kernel-Definition hinzufügen und bieten eine geeignete Definition für AVX_VECTOR_SIZE (die nicht definiert zu sein scheint In Ihrem Beispiel AFAICT), kompiliert und läuft Ihr Code korrekt für mich.

Answer 2

ich gefunden habe, könnte es möglich sein, das Problem ohne explizite Template-Instantiierung zu beheben:

class VectorDevice 
{ 
    T* data; 
    size_t size; 
    int dim, dim_aligned; 

    struct A 
    { 
     __host__ 
     A() 
     { 
      bool neverCalled = true; 
      if (!neverCalled) 
      { 
       deviceSizeOf<T><<<1, 1>>>(NULL); 
       CUDA_ERR_CHECK(cudaGetLastError()); 
       CUDA_ERR_CHECK(cudaDeviceSynchronize()); 
      } 
     } 
    } a; 

public : 

    __host__ __device__ 
    VectorDevice() : data(NULL), size(0), dim(0) { } 

    #pragma hd_warning_disable \ 
    #pragma nv_exec_check_disable 
    __host__ __device__ 
    VectorDevice(int dim_) : data(NULL), size(0), dim(dim_) 
    { 
     ... 
    } 

    ... 
}; 

...