我想定义一个在设备和主机代码之间共享的指针类型,并在内部将设备和主机指针存储到共享内存。我希望它在编译时确定实际返回的指针:
#define F inline __host__ __device__
class SharedMemory;
/**
*
* Can only be set by allocating shared memory.
*/
template<typename T>
class SharedMemoryPtr {
public:
SharedMemoryPtr() : hptr(0), dptr(0) {}
//F ~ SharedMemoryPtr() {cudaFreeHost(hptr);} // Should be freed explicitly (?)
// TODO: Don't allow copying/overwriting (at least not without freeing memory...)
F T& operator() () {
#ifdef __CUDACC__
return *dptr;
#else
return *hptr;
#endif
};
F T* operator-> () {
#ifdef __CUDACC__
return dptr;
#else
return hptr;
#endif
};
F T& operator * () {
#ifdef __CUDACC__
return *dptr;
#else
return *hptr;
#endif
};
F T& operator[] (__int64 i) const {
#ifdef __CUDACC__
return *(dptr + i);
#else
return *(hptr + i);
#endif
};
friend SharedMemory;
// TODO: Make obsolete (possible?)
T * getHptr() {return hptr;}
T * getDptr() {return dptr;}
private:
T *hptr, *dptr;
};
class SharedMemory {
public:
template<typename T>
static SharedMemoryPtr<T> allocate(int count = 1) {
assert(count > 0);
SharedMemoryPtr<T> sptr;
cutilSafeCall(
cudaHostAlloc(&sptr.hptr, sizeof(T) * count, cudaHostAllocMapped));
assert(sptr.hptr);
cutilSafeCall(
cudaHostGetDevicePointer(&sptr.dptr, sptr.hptr, 0));
assert(sptr.dptr);
return sptr;
}
};
只要我在 cpp 文件(__CUDACC__
从未定义过)或 .h 文件(__CUDACC__
仅当函数被 cu 文件中的某个函数使用时才定义)中的代码中使用这些指针,这就可以正常工作。但是在 .cu 文件中的一个__host__
函数中,我得到了 .cu 文件devptr
。显然,.cu 文件由 nvcc 专门处理。是否有其他一些仅为__global__
和__device__
函数定义的预处理器宏,而不仅仅是 nvcc 碰巧处理的所有内容?还是我需要分开我的代码?