用户
 找回密码
 立即注册
xinliangyu 该用户已被删除
发表于 2013-10-25 22:11:15
101616
先要谢谢版主前面的解答,由于本人初学cuda,并且懒惰导致忽视了版主的辛苦解答!
现在遇到新的问题,不知道有人帮我解答下不?
我现在将204张512*512张图像绑定到三维纹理中(没有偏移设置),现在想要按照volumeRender的方法在光线方向设置51个步长(grid(32,32),block(16,16)),每次tex3D提取一次数值保存,形成一个512*512*51张图像输出,但是现在遇到的图像输出有误,代码附上:
        const int maxSteps = 51;//Maxsteps和tstep的设置查了下,好像没有太多中文解释,不知道这里有错误没
        const float tstep = 0.02f;
        const float3 boxMin = make_float3(-1.0f, -1.0f, -1.0f);
        const float3 boxMax = make_float3(1.0f, 1.0f, 1.0f);

        uint x = blockIdx.x*blockDim.x + threadIdx.x;
        uint y = blockIdx.y*blockDim.y + threadIdx.y;

        if ((x >= imageW) || (y >= imageH)) return;

        float u = (x / (float) imageW)*2.0f-1.0f;
        float v = (y / (float) imageH)*2.0f-1.0f;

        // calculate eye ray in world space
        //执行了下面语句之后,eyeRay.o ={0, 0, 4}
        Ray eyeRay;
        eyeRay.o = make_float3(mul(c_invViewMatrix, make_float4(0.0f, 0.0f, 0.0f, 1.0f)));
        eyeRay.d = normalize(make_float3(u, v, -2.0f));
        eyeRay.d = mul(c_invViewMatrix, eyeRay.d);

        // find intersection with box
        float tnear, tfar;
        int hit = intersectBox(eyeRay, boxMin, boxMax, &tnear, &tfar);

        if (!hit) return;
        if (tnear < 0.0f) tnear = 0.0f;     // clamp to near plane

        float t = tnear;
        float3 pos = eyeRay.o + eyeRay.d*tnear;
        float3 step = eyeRay.d*tstep;

        for (int i=0; i<maxSteps; i++)
        {
                // read from 3D texture
                // remap position to [0, 1] coordinates
                //原始数据uchar已经被转成float型了,位置也是float型,从float型位置再读图像原始值
                //pos.x = pos.x*volumeSize.width/2.0f+volumeSize.width/2.0f;        // map position.x to [0, 512] coordinates
                //pos.y = pos.y*volumeSize.height/2.0f+volumeSize.height/2.0f;      // map position.y to [0, 512] coordinates
                //pos.z = pos.z*volumeSize.depth/2.0f+volumeSize.depth/2.0f;        // map position.z to [0, 331] coordinates
                float sample = tex3D(tex, pos.x*0.5f+0.5f, pos.y*0.5f+0.5f, pos.z*0.5f+0.5f);

                t += tstep;
                if (t > tfar) break;

                pos += step;
                outputData[y*imageW + x + i*imageW*imageH] =  (short)round(sample);
        }

使用道具 举报 回复
发表于 2013-10-25 22:28:45
楼主您好,

您都承认您再复制的时候没有指定偏移量,您怎可能读取出来正确的值?您这样后续的204次复制将覆盖前一次的。

以及,如果要解决问题,请重新阅读上次的解决方案。不要无视我们的建议。
(以及,您这次连cuda array的建立,数据的传输,Texture reference是如何声明的,都没有。您让论坛如何为您解决?这不是为难人么。)

请您三思。
感谢深夜来访。
使用道具 举报 回复 支持 反对
发表于 2013-10-26 09:57:34
玫瑰幻想 发表于 2013-10-25 22:28
楼主您好,

您都承认您再复制的时候没有指定偏移量,您怎可能读取出来正确的值?您这样后续的204次复制将 ...

谢谢版主的解答,我看到volumeRender的数据由指针绑定纹理时就没有设置偏移量,不过它由指针直接只想fread进来的数据,关于texture reference的设置我都参考了volumerender中(归一化),然后它就可以在沿着ray叠加各个(pos+step)上的像素值,然后在与颜色纹理相乘得到最终的数据。
还有我真的不是有意忽视版主认真的解答,实在是由于本人出学,只能参考samples种的东西来做的;
那我把数据绑定那块代码也附上吧。
数据读取
short *loadRawFile(char *filename, size_t size)

{

        FILE *fp = fopen(filename, "rb");

        if (!fp)

        {

                fprintf(stderr, "Error opening file '%s'\n", filename);

                return 0;

        }

       

        short *data = (short*)malloc(size);

        if (NULL == data)

        {

                printf("Malloc data failue!\n");

                return 0;

        }

       

        size_t read = fread(data, 2, 512*512*408, fp);

        if (0 == read)

        {

                printf("Read data failue!\n");

                return 0;

        }

        fclose(fp);

        //test the file data

        //ofstream out("out.dat");

        //int i = 512*512*1;

        //int j = 0;

        //while (j<512*512)

        //{

        //        //printf("%d\n",dat[i++]);

        //        out<<data[i++]<<" ";

        //        if((i+1)/512 == 0)

        //                out<<endl;

        //        j++;

        //}

        //out.close();

        printf("Read '%s', %d bytes\n", filename, read);

        return data;

}
把数据的按照204张循环拷贝给cudaArray
cudaExtent volumeSize = make_cudaExtent(512, 512, 204);

        size_t VolumeSize = volumeSize.width*volumeSize.height* 408 *sizeof(VolumeType);

        size_t newVolumeSize = volumeSize.width*volumeSize.height* 204 *sizeof(VolumeType);

        size_t subVolumeSize = volumeSize.width*volumeSize.height* 51 *sizeof(VolumeType);

        size_t PicSize = volumeSize.width * volumeSize.height *sizeof(VolumeType);

       

        short   *h_volume = loadRawFile(path, VolumeSize);

        if ( NULL == h_volume)

        {

                printf("h_volume is NULL!\n");

                return ;

        }



        short  *out_volume = (short*)malloc(512*512*102*sizeof(short)); //save data after Interpolation

        if ( NULL == out_volume)

        {

                printf("h_volume is NULL!\n");

                return ;

        }

       

        short   *h_Involume = (short*) malloc(512*512*204*sizeof(short));//¿œ±ŽžøÎÆÀíÊýŸÝ

        if ( NULL == h_Involume)

        {

                printf("h_volume is NULL!\n");

                return ;

        }



        int    iTimes = 1;//408/102 ;//cuda calculate times

        int    iTime = 0;

        while(iTime < iTimes)

        {               

                memcpy(h_Involume, h_volume+512*512*204 , 512*512*204*sizeof(short));



                ofstream test1("myOut1.dat");

                int ii1=512*512*0;

                while (ii1<512*512*1)

                {

                test1<<h_Involume[ii1++]<<" ";

                }

                test1.close();

                ofstream test2("myOut2.dat");

                int ii2=512*512*1;

                while (ii2<512*512*2)

                {

                test2<<h_Involume[ii2++]<<" ";

                }

                test2.close();

                ofstream test3("myOut3.dat");

                int ii3=512*512*2;

                while (ii3<512*512*3)

                {

                test3<<h_Involume[ii3++]<<" ";

                }

                test3.close();

                ofstream test4("myOut4.dat");

                int ii4=512*512*3;

                while (ii4<512*512*4)

                {

                test4<<h_Involume[ii4++]<<" ";

                }

                test4.close();



                cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<VolumeType>();

                checkCudaErrors(cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumeSize));

                // copy data to 3D array

                cudaMemcpy3DParms copyParams = {0};

                copyParams.srcPtr   = make_cudaPitchedPtr(h_Involume, volumeSize.width*sizeof(VolumeType), volumeSize.width, volumeSize.height);

                copyParams.dstArray = d_volumeArray;

                copyParams.extent   = volumeSize;

                copyParams.kind     = cudaMemcpyHostToDevice;

                checkCudaErrors(cudaMemcpy3D(&copyParams));

               

                // set texture parameters

                tex.normalized = true;                      // access with normalized texture coordinates

                tex.filterMode = cudaFilterModeLinear;      // linear interpolation

                tex.addressMode[0] = cudaAddressModeClamp;  // clamp texture coordinates

                tex.addressMode[1] = cudaAddressModeClamp;

                // bind array to 3D texture

                checkCudaErrors(cudaBindTextureToArray(tex, d_volumeArray, channelDesc));

        cudaFreeArray(d_volumeArray);

                getLastCudaError("Kernel execution failed");



                short* dOutData;

                cudaMalloc((short **)&dOutData,512*512*51*sizeof(short));



                dim3 dimBlock(32, 32, 1);

                dim3 dimGrid(16, 16, 1);

                // Warmup

                transformKernel<<<dimGrid, dimBlock>>>(dOutData, width, height);

                getLastCudaError("Kernel execution failed");



                checkCudaErrors(cudaDeviceSynchronize());

                StopWatchInterface *timer = NULL;

                sdkCreateTimer(&timer);

                sdkStartTimer(&timer);



                // Execute the kernel

                transformKernel<<<dimGrid, dimBlock>>>(dOutData, width, height);

                // Check if kernel execution generated an error

                getLastCudaError("Kernel execution failed");



                cudaMemcpy(out_volume + 512*512*51*iTime, dOutData, (512*512*51*sizeof(short)), cudaMemcpyDeviceToHost);

                ofstream out("fisrtPage.dat");

                int j=0;

                while (j<512*512)

                {

                        out<<out_volume[j]<<endl;

                        j++;

                }

                out.close();

                cudaFree(dOutData);

                iTime++;

        }

        free(h_Involume);

        free(h_volume);
kernel中函数
__global__ void transformKernel(short *outputData,int imageW, int imageH)

{

        const int maxSteps = 10;

        const float tstep = 0.02f;

        const float3 boxMin = make_float3(-1.0f, -1.0f, -1.0f);

        const float3 boxMax = make_float3(1.0f, 1.0f, 1.0f);



        uint x = blockIdx.x*blockDim.x + threadIdx.x;

        uint y = blockIdx.y*blockDim.y + threadIdx.y;



        if ((x >= imageW) || (y >= imageH)) return;



        float u = (x / (float) imageW)*2.0f-1.0f;

        float v = (y / (float) imageH)*2.0f-1.0f;



        // calculate eye ray in world space

        //ÖŽÐÐÁËÏÂÃæÓïŸäÖ®ºó,eyeRay.o ={0, 0, 4}

        Ray eyeRay;

        eyeRay.o = make_float3(mul(c_invViewMatrix, make_float4(0.0f, 0.0f, 0.0f, 1.0f)));

        eyeRay.d = normalize(make_float3(u, v, -2.0f));

        eyeRay.d = mul(c_invViewMatrix, eyeRay.d);



        // find intersection with box

        float tnear, tfar;

        int hit = intersectBox(eyeRay, boxMin, boxMax, &tnear, &tfar);



        if (!hit) return;

        if (tnear < 0.0f) tnear = 0.0f;     // clamp to near plane



        float t = tnear;

        float3 pos = eyeRay.o + eyeRay.d*tnear;

        float3 step = eyeRay.d*tstep;



        for (int i=0; i<maxSteps; i++)

        {

                // read from 3D texture

                // remap position to [0, 1] coordinates

                //Ô­ÊŒÊýŸÝucharÒÑŸ­±»×ª³ÉfloatÐÍÁË£¬Î»ÖÃÒ²ÊÇfloatÐÍ£¬ŽÓfloatÐÍλÖÃÔÙ¶ÁÍŒÏñԭʌֵ

                //pos.x = pos.x*volumeSize.width/2.0f+volumeSize.width/2.0f;        // map position.x to [0, 512] coordinates

                //pos.y = pos.y*volumeSize.height/2.0f+volumeSize.height/2.0f;      // map position.y to [0, 512] coordinates

                //pos.z = pos.z*volumeSize.depth/2.0f+volumeSize.depth/2.0f;        // map position.z to [0, 331] coordinates

                float sample = tex3D(tex, pos.x*0.5f+0.5f, pos.y*0.5f+0.5f, pos.z*0.5f+0.5f);



                t += tstep;

                if (t > tfar) break;



                pos += step;

                outputData[y*imageW + x + i*imageW*imageH] =  (short)round(sample);

        }

}

再次真心谢谢版主!麻烦了
使用道具 举报 回复 支持 反对
发表于 2013-10-26 19:03:31
本帖最后由 横扫千军 于 2013-10-26 19:04 编辑
xinliangyu 发表于 2013-10-26 09:57
谢谢版主的解答,我看到volumeRender的数据由指针绑定纹理时就没有设置偏移量,不过它由指针直接只想frea ...

楼主,

虽然现在是周末,但也请认真发帖。
抄袭要抄认真,提问要给要求的信息(例如texture reference的定义你都不发!都反复要求了N遍,横跨N个主题贴了,要求你发N次始终无视,你以为我是神仙能自动看穿你的硬盘里的内容么?)。

你上来一堆代码,
这种malloc(make_cudaExtent(512, 512, 204));
这种memcpy(h_Involume, h_volume+512*512*204 , 512*512*204*sizeof(short));

你清醒你在干什么么?

你我如果交换身体,你会怎么做?三思吧!!
楼主自重!
使用道具 举报 回复 支持 反对
发表于 2013-10-26 22:08:04
横扫千军 发表于 2013-10-26 19:03
楼主,

虽然现在是周末,但也请认真发帖。

呵呵 。。。。。。。。。。。。
使用道具 举报 回复 支持 反对
发表于 2013-10-26 22:09:48
如果我懂我来该版做什么。。。
使用道具 举报 回复 支持 反对
发表于 2013-10-26 23:13:05
本帖最后由 xinliangyu 于 2013-10-26 23:45 编辑
横扫千军 发表于 2013-10-26 19:03
楼主,

虽然现在是周末,但也请认真发帖。

首先还是感谢版主前面的解答,但是我实在觉得不舒服,想跟版主把话说清楚。
我前面说过我参考的volumeRender的实例,在附上代码时忽略了
texture<short, 3, cudaReadModeNormalizedFloat> tex;  //这一行
但是版主说要求了N次的texture reference在volumeRender中就没有用到,并且我也没有有意忽视你的意见。
并且我怎么就抄袭了?版主是怎么定义抄袭的?在你学习时是怎么进步的?
不知道版主有没有听过 “授业有先后,术业有专攻”?
还有我附上的代码上哪有malloc(make_cudaExtent(512, 512, 204));这一行?
并且版主鄙视我的这行代码?memcpy(h_Involume, h_volume+512*512*204 , 512*512*204*sizeof(short));
这里哪有错误?我要读进来的照片内存部分赋给h_Involume中,有什么错误?,难道是因为在.cu文件?可是我都测试过读进来内存,经过保存显示照片可以这么做的!

最后版主要如果交换身体来换位思考,那楼主有没有想过我为什么周末发帖。
希望版主能够给出上面的解释,再次谢谢。。。
使用道具 举报 回复 支持 反对
发新帖
您需要登录后才可以回帖 登录 | 立即注册