CUDA结构体内指针问题

初学CUDA,想请教一下:
/** a node in a k-d tree */
#pragma pack(16)
struct kd_node
{
int index;
int index_left;
int index_right;
int ki; /< partition key index */
double kv; /
< partition key value */
int leaf; /< 1 if node is a leaf, 0 otherwise /
struct feature
features; /
< features at this node */
int n; /< number of features /
struct kd_node
kd_left; /
< left child /
struct kd_node
kd_right; /**< right child /
};
#pragma pack()
struct feature
{
int d; /descriptor length/
double descr[FEATURE_MAX_D]; /descriptor/
void
feature_data; /user-definable data/
};

然后定义两个结构体指针:
struct kd_node* d_expl;
cudaMalloc((void**)&d_expl,n2*sizeof(struct kd_node));

struct feature* d_tree_feat;
cudaMalloc((void**)&d_tree_feat,n2*sizeof(struct feature));

在global函数中这样写:
int tid=threadIdx.x+blockIdx.x*blockDim.x;
d_tree_feat[tid]=d_expl[tid].features[0];
貌似不对,我想是不是因为没有给d_expl中的feature分配显存造成的,要是这样的话,这个应该怎样分配显存啊???
有没有谁能给解答一下,急。。。这个问题纠缠了一个一星期了,没时间浪费了。先谢谢了。

补充一下,struct kd_node中的int n,为kd_node中features的个数,对不同的kd_node,features的个数是不同的。

请看此文。此文的3#中,ICE写了一千五百字。可以说是这种问题的通用解答了。

http://cudazone.nvidia.cn/forum/forum.php?mod=viewthread&tid=6820&extra=page%3D1

请仔细阅读ICE的给出的解决方案。

BZ,我是初学者,就是看那个没太明白才问的,能不能就我这个说一下代码怎么写。不需要太多字,就是我这个代码怎么改一下。谢谢。