指针数组

加速计算 CUDA
1

在主机端定义全局指针数组float * dev_MatFeat[30];

//第一个函数使用，然后循环初始化
fun1（int num…）
{
cudaMalloc((void **)&dev_MatFeat[num],out0out1324);
cudaMemset(dev_MatFeat[num],0,out0out1324);
histnorm2feat<<<out1, out0>>>(dev_hist,dev_norm,dev_MatFeat[num]);
cudaMemcpy(src[num],dev_MatFeat[num],out0out132*4,cudaMemcpyDeviceToHost);
}

//第二次函数使用，同一个cu文件fun2（int t）
{

// cudaMemcpy(dev_MatFeat[t] ,src[t] ,out0out132*4 ,cudaMemcpyHostToDevice);
padarray<<< blockFeat, src.feat[t].size[0]>>>(dev_MatFeat[t],dev_MatDst[t],padx,pady);
//使用注释和没注释得到不同的结果，不知道为什么

}

2

LZ您好：

1：根据您给出的代码，您的操作都是在默认流里面的，他们是顺序执行的，原则上后面kernel启动的时候，前面的kernel已经结束了。

2：根据您注释掉的代码，第一个循环中注释掉了dev_MatFeat的回拷，第二段中注释掉了重新申请空间，和将前面回拷内容copy进显存的过程，以及注释掉了一个kernel。

值得注意的是，您注释掉的kernel有一个参数是dev_MatDst，如果该kernel生成了dev_MatDst的内容，你注释掉这个kernel，那么dev_MatDst的值就是不正确的，此时您最后一个kernel使用dev_MatDst作为参数，可能就因此获得了错误的结果。

您的代码大致分析如此，供您参考。

祝您一切顺利~

3

第二段代码这样 表述可能好些，
我重新修改了代码了，看第一楼，就是架上来回复制部分和没加上得到不同的结果,还有我第一段代码的哪个num不是按顺序的，而是按 5 ，10，15，等打乱的顺序开辟指针的，没有重复开辟，从num的值从0-23都是一次，第二段代码的t是按顺序，从0一直到23

4

这样使用指针数组对不对呢？
我内核函数式这样定义的
global void padarray(float src,float dst,int padx,int pady);
如果使用没有问题的话我检查其他方面的原因

5

LZ你好：

1：建议您以后修改代码请在回复楼中重新完整贴出，而不是去修改原帖。您直接修改了1#的代码，使得我2#当时的回答无法和您修改后的代码对应，这将造成其他看帖的人十分迷惑，以及我也无法回溯当时的情况。

2：请将您的代码补全写在后面的回复楼层中，并且将您3#的叙述用实际代码表达出来，以免造成理解上的差异。我看您的描述并无法明白您的行为和意图，也许您自己是明白的，这个正常，因为您有一系列关于您问题的先验知识以及您明白您的具体问题，但是我并无法凭空得知。

3：以及，您可能还需要检查下没贴出来的部分，是否有影响。

大致如此，祝您好运~

6

LZ您好：

您kernel函数的参数类型是float*；您调用函数时，实参的类型与之匹配即可。
dev_MatFeat是指针数组，每个元素是float*类型的，是匹配的。
dev_MatDst您没有给出定义，无法判定。

大致如此，祝您好运~

7

void calcfeatpyramid_fast(…)
{

for (int i=0;i<5;i++) //interval =5
{

cuda_features_4bin(feat[i],feat[i+interval],i,i+5);

for (int j=i+5;j<19;j=j+5)
{
cuda_features_8bin(feat[j+interval],j+interval);
}

}
}

extern “C” void cuda_features_4bin(_feat& feat,_feat& feat8bin,int num4,int num8)
{

…
feat.ft=new float[out0out132];
memset(feat.ft,0,out0out132*4);

cudaMalloc((void **)&dev_MatFeat[num4],out0out1324);
cudaMemset(dev_MatFeat[num4],0,out0out1324);

histnorm2feat<<<out1, out0>>>(dev_hist,dev_norm,dev_MatFeat[num4]);
cudaMemcpy(feat.ft,dev_MatFeat[num4],out0out132*4,cudaMemcpyDeviceToHost);

…

feat8bin.ft=new float[out0out132];
memset(feat8bin.ft,0,out0out132*4);

cudaMalloc((void **)&dev_MatFeat[num8],out0out1324);
cudaMemset(dev_MatFeat[num8],0,out0out1324);

histnorm2feat<<<out1, out0>>>(dev_hist8bins,dev_norm8bins,dev_MatFeat[num8]);
cudaMemcpy(feat8bin.ft,dev_MatFeat[num8],out0out132*4,cudaMemcpyDeviceToHost);

}

extern “C” void cuda_features_8bin(_feat& feat,int num)
{

feat.ft=new float[out0out132];
memset(feat.ft,0,out0out132*4);

cudaMalloc((void **)&dev_MatFeat[num],out0out1324);
cudaMemset(dev_MatFeat[num],0,out0out1324);

histnorm2feat<<<out1, out0>>>(dev_hist,dev_norm,dev_MatFeat[num]); //0.6

cudaMemcpy(feat.ft,dev_MatFeat[num],out0out132*4,cudaMemcpyDeviceToHost);

}

extern “C” void cuda_padarray(…)
{

for (int t=0;t<24;t++)
{

dstfeat[t].ft=new float[dstfeat[t].size[0]*dstfeat[t].size[1]*dstfeat[t].size[2]];
memset(dst.feat[t].ft,0,sizeof(float)*dst.feat[t].size[0]*dst.feat[t].size[1]*dst.feat[t].size[2]);

cudaMalloc((void **)&dev_MatDst[t],dstfeat[t].size[0]*dstfeat[t].size[1]*dstfeat[t].size[2]*4);
cudaMemset(dev_MatDst[t],0,dstfeat[t].size[0]*dstfeat[t].size[1]*dstfeat[t].size[2]*4);

// cudaMemcpy(dev_MatFeat[t],src.feat[t].ft,src.feat[t].size[0]*src.feat[t].size[1]*src.feat[t].size[2]*4,
//cudaMemcpyHostToDevice);

padarray<<< blockFeat, dst.feat[t].size[0]>>>(dev_MatFeat[t],dev_MatDst[t],padx,pady);

cudaMemcpy(dstfeat[t].ft,dev_MatDst[t],dstfeat[t].size[0]*dstfeat[t].size[1]*dstfeat[t].size[2]*4,cudaMemcpyDeviceToHost);

cudaFree(dev_MatFeat[t]);

}

}
好吧，我把主要代码都贴出来
float * dev_MatFeat[30];
float * dev_MatDst[30];
这两个是全局变量
注释部分去掉和没去掉得到不同的结果。要是觉得代码太长就算了[/i]srcfeat是前面三个求得的结果feat

8

LZ您好，感谢您贴出代码，这样比较清楚。

根据您提供的信息，dev_MatDst也是指针数组，其每个元素的类型也是合适的。

您的代码中对dev_MatFeat每个元素（指针）申请了空间，启动kernel生成了内容，并回拷了数据。这个虽然不是顺序从第一个元素到最后一个元素，但并无问题，既无重复，也无遗漏。

之后，您在后面的代码中顺序使用了dev_MatFeat各个元素，这个也是没有问题的。

原则上您申请的空间，的确可以反复使用的，但是您却观察到了直接使用和host端copy过来使用，结果不一致。
考虑到您是每次kernel生成对应数据写入显存之后，立即回拷到host的，您host的数据可应该是无误的。
那么，您可以检查一下，device端这些对应的数据是否被人为改写过，或者检查下是否因为越界而意外改变了这些数据。

目前暂时想到这个可能，您或许需要仔细检查一下了。

大致建议如上，祝您调试顺利~