列出帖子

该操作将允许你查看该会员所有的帖子,注意你只能看到你有权限看到的板块的帖子。


只显示主题 - 屠戮人神

页: [1]
1
CUDA / 请问当cuda中出现unknown error问题怎么解决?
« 于: 三月 02, 2019, 10:12:03 pm »
当程序数据小的时候正常运行,数据大时就会跳过Kernel不运行核函数
用下面这段小程序测试错误出现  unknown error,请问该如何解决??? 


error_check = cudaGetLastError();
        if (error_check != cudaSuccess) {
                printf("%s\n", cudaGetErrorString(error_check));
                system("pause");
               
        }

源程序大致如下:
void setKernel (cuComplex *d_Z,……)
{      int idx= threadIdx.x + blockIdx.x*blockDim.x;
        int idy= threadIdx.y + blockIdx.y*blockDim.y;
        for (int i=idx;i < a; i = i + blockDim.x * gridDim.x)
        {
            …………
            for (int j=idy;j < b; j = j + blockDim.y * gridDim.y)
           {
             …………
            d_Z[i*b+j]=……
           }//-------j
       }  //-------i 

}
void set (h_Z)
{       
        cuComplex *h_Z=new cuComplex[a*b];
        cudaMalloc((void**)&d_Z, a*b* sizeof(cuComplex));

        const dim3 blockSize(8, 8);
        const dim3 gridSize(32, 32);
        SetKernel << <gridSize, blockSize >> > (d_Z……);

        cudaMemcpy(h_Z, d_Z, a*b* sizeof(cuComplex), cudaMemcpyDeviceToHost);

        delete[]h_Z;
        cudaFree(d_Z);
}


2
CUDA / 核函数拆分和依赖问题
« 于: 三月 05, 2019, 05:31:22 pm »
是这样,我这边对大量的数据进行GPU运算,之前是放在同一个核函数处理的,主要处理逻辑分为两步。但发现有大量的数据不满足第二步的条件,这样其实有很大一部分GPU资源是浪费掉的,所以考虑对核函数进行拆分。由于第二步的运算依赖于第一步运算计算结果,且第一部运算的结果都是散列的,个数不固定,需要进行整理和统计后,才能进行第二步的计算。请问大家有没有好的思路和策略能够高效的对核函数进行拆分,实现性能提升?

3
处理同样大小的数据时,分别使用thrust::reduce和自己定义reduce函数,经测得自己定义的速度会快一点.我想请教一下,什么时候适合调用thrust库内的函数?

4
环境 win10+32G内存。需要的显存超过的GPU显存。只能尝试着用零拷贝内存。

1、cuda自带的例子 simpleZeroCopy,对于32G的内存,只能分配大概14.25G锁页内存,超过这个容量,cudaHostAlloc就失败了,但是明明还有10+G的空闲内存啊,这个失败的原因是什么?在cuda社区查看,似乎linux/windows+专业卡是可以几乎分配所有的内存的,这个是window驱动的问题吗?有没有谁能够指点一下?谢谢。

2、另外,window下cudaMallocManaged 什么时候能够支持GPU Memory Oversubscription啊,看文档Linux下是可以的,有谁用过吗?

谢谢

5
CUDA / 关于内核并发
« 于: 二月 26, 2019, 03:49:22 pm »
之前发的那个帖子好像被删掉了不知道怎么回事,是这样的,我的平台是win10 驱动WDDM2.1,卡是gtx1080(这个配置是会禁用Hyper-Q吗?)
我使用流处理想要实现内核并发,大概过程是这样的
         data1.upload(host_src_pl1, stream0);
         Kernel1(ImageShiftGpu1, layer, stream0);
         data2.upload(host_src2_pl1, stream1);
         Kernel2(ImageShiftGpu2, layer, stream1);
         data3.upload(host_src3_pl1, stream2);
         Kernel3(ImageShiftGpu3, layer, stream2);

测试了一下,他可以屏蔽掉data2,data3的上传的时间,但是三个内核并没有并行运行,就是我以为它应该按照图中的左边执行的,但实际上它可能是按照图里的右边执行的,大概是因为什么原因呢?


6
配置环境是win10+VS2015+CUDA 10.0+Titan RTX。使用cuda 8和9.0进行编译运行时,一切正常,切换到CUDA 10.0时,出现类似于:

c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\channel_descriptor.h(104): error C2894: 模板不能声明为有“C”链接
1>c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\channel_descriptor.h(104): error C2733: “cudaCreateChannelDesc”: 不允许重载函数的第二个 C 链接
1>  c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\cuda_runtime_api.h(7398): note: 参见“cudaCreateChannelDesc”的声明
1>c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\channel_descriptor.h(137): error C2894: 模板不能声明为有“C”链接
1>c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\channel_descriptor.h(137): error C2733: “cudaCreateChannelDesc”: 不允许重载函数的第二个 C 链接
1>  c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\channel_descriptor.h(104): note: 参见“cudaCreateChannelDesc”的声明
……

莫名奇妙的错误,可能原因是什么?

7
CUDA / 使用cuda对大量数据进行排序时出错
« 于: 二月 28, 2019, 11:48:19 am »
思路:输入数据的基数排序,通过扫描操作实现按位与后0和1的顺序排列
错误:96*256个数据时,排序操作是正确的,而97*256个数据时排序结果不正确,经排查出错操作为包含扫描操作

疑问:cuda进行包含扫描操作对数据大小是有什么限制吗?为什么超出96*256扫描就不对了?

*****************显卡属性***********
使用GPU device 0: GeForce GTX 1070
SM的数量:15
每个线程块的共享内存大小:48 KB
每个线程块的最大线程数:1024
每个SM的最大线程数:2048
每个SM的最大warp数:64

程序代码: [选择]
#include <stdio.h>
#include <iostream>
#include <cuda_runtime.h>
#include <thrust/reduce.h>
#include <thrust/device_vector.h>

using namespace std;
__global__ void warmup(){}
__global__ void  cdf(unsigned int *cdf,unsigned int *in,const int rows,const int cols,const unsigned int i)
{
    int c=blockIdx.x*blockDim.x+threadIdx.x;
    int r=blockIdx.y*blockDim.y+threadIdx.y;
    int idx=r*cols+c;
    if(c>=cols||r>=rows)return;
    unsigned int mask=1<<i;
    unsigned int histo[2]={1,0};
    unsigned int bin = (in[idx] & mask) >> i;
    cdf[idx]=histo[bin];
}

__global__ void  cdf_1(unsigned int *cdf,unsigned int *in,const int rows,const int cols,const unsigned int i)
{
    int c=blockIdx.x*blockDim.x+threadIdx.x;
    int r=blockIdx.y*blockDim.y+threadIdx.y;
    int idx=r*cols+c;
    if(c>=cols||r>=rows)return;
    unsigned int mask=1<<i;
    unsigned int bin = (in[idx] & mask) >> i;
    cdf[idx]=bin;
}

__global__ void ex_scan(unsigned int *out,unsigned int *tmp,const int i,const int rows, const int cols)
{
    int c=blockIdx.x*blockDim.x+threadIdx.x;
    int r=blockIdx.y*blockDim.y+threadIdx.y;
    int idx=r*cols+c;
    if(c>=cols||r>=rows)return;
    int k;
    k=1<<i;
        out[idx]=tmp[idx];
    if(idx-k<0)return;
    unsigned  int b=0;
    b=tmp[idx-k];
    out[idx]+=b;
    __syncthreads();
    tmp[idx]=out[idx];

}

__global__ void sort(unsigned int *d_vals,unsigned int *d_scan,unsigned int *d_cdf,unsigned int *d_in,const int rows,const int cols,const int index)
{
    int c=blockIdx.x*blockDim.x+threadIdx.x;
    int r=blockIdx.y*blockDim.y+threadIdx.y;
    int idx=r*cols+c;
    if(c>=cols||r>=rows)return;
    if(d_cdf[idx]==0)return;
    if(idx==20)printf("d_cdf[idx]:%d\t,d_scan[idx]:%d\t,d_in[idx]:%d\t\n",d_cdf[idx],d_scan[idx],d_in[idx]);
    d_vals[index+d_scan[idx]-1]=d_in[idx];
}

void func_cuda()
{
    warmup<<<1,1>>>();
    unsigned int *h_in,*d_in,*d_vals,*d_scan,*d_tmp,*h_scan,*h_vals;
    int rows,cols;
    rows=96;
    cols=256;
    int bytes,numElems,num;

    numElems=rows*cols;
    num=(int)log2f(numElems)+1;
    bytes=rows*cols*sizeof(unsigned int);
    h_in = (unsigned int *) malloc (bytes);
    h_vals = (unsigned int *) malloc (bytes);
    for(unsigned int i=0;i<rows;i++)
    {
        for(unsigned int j=0;j<cols;j++)
        {
            h_in[i*cols+j]=numElems-(i*cols+j);
        }
    }

    h_scan = (unsigned int *) malloc (bytes);

    thrust :: device_vector <unsigned int> d_cdf(numElems);
    cudaMalloc(&d_in,bytes);
    cudaMalloc(&d_vals,bytes);
    cudaMalloc(&d_scan,bytes);
    cudaMalloc(&d_tmp,bytes);
    cudaMemcpy(d_in,h_in,bytes,cudaMemcpyHostToDevice);

    dim3 blocks((cols-1+32)/32,(rows-1+32)/32,1);
    dim3 threads(32,32,1);

    for (unsigned int i = 0; i <8*sizeof(unsigned int) ; i++)
    {
        cdf<<<blocks,threads>>>(thrust::raw_pointer_cast(d_cdf.data()),d_in,rows,cols,i);
        int sum = thrust :: reduce (d_cdf.begin () , d_cdf.end () , ( int ) 0, thrust :: plus <int >());
        cudaMemcpy(d_tmp,thrust::raw_pointer_cast(d_cdf.data()),bytes,cudaMemcpyDeviceToDevice);
        for(int j=0;j<num;j++)
        {
        ex_scan<<<blocks,threads>>>(d_scan,d_tmp,j,rows, cols);
        }
        sort<<<blocks,threads>>>(d_vals,d_scan,thrust::raw_pointer_cast(d_cdf.data()),d_in,rows,cols,0);

        cdf_1<<<blocks,threads>>>(thrust::raw_pointer_cast(d_cdf.data()),d_in,rows,cols,i);
        cudaMemcpy(d_tmp,thrust::raw_pointer_cast(d_cdf.data()),bytes,cudaMemcpyDeviceToDevice);
        for(int j=0;j<num;j++)
        {
        ex_scan<<<blocks,threads>>>(d_scan,d_tmp,j,rows, cols);
        }
        sort<<<blocks,threads>>>(d_vals,d_scan,thrust::raw_pointer_cast(d_cdf.data()),d_in,rows,cols,sum);

        cudaMemcpy(d_in,d_vals,bytes,cudaMemcpyDeviceToDevice);

    }
    cudaMemcpy(h_vals,d_vals,bytes,cudaMemcpyDeviceToHost);

    for(int j=0;j<rows;j++)
    {
        for(int k=0;k<cols;k++)
        {
            cout<<h_vals[j*cols+k]<<"   ";
        }
        cout<<endl;
    }

    free(h_in);
    free(h_vals);
    free(h_scan);//释放数组
    cudaFree(d_in);
    cudaFree(d_vals);
    cudaFree(d_scan);
    cudaFree(d_tmp);

8
CUDA / 我有许多kernel函数
« 于: 二月 26, 2019, 02:40:03 pm »
我有许多kernel函数,我单独的执行每个kernel,统计他们的时间的和记为t1。然后我现在把这些kernel放到一起执行,他们之间有数据依赖问题(就是一些kernel的输出是另外一些kernel的输入),已经通过图来描述清楚了,这种方式的时间和为t2。发现t2比t1大很多。另外用nvprof测得每个kernel运行的时间差异也很大,第二种方式是第一种方式的10倍的运算时间。请教下这可能得原因是什么呢,谢谢

9
本帖最后由 pickou 于 2019-2-26 16:57 编辑

如题,我有许多kernel,我单独的执行每个kernel,统计他们的时间的和记为t1。然后我现在把这些kernel放到一起执行,他们之间有数据依赖问题(就是一些kernel的输出是另外一些kernel的输入),已经通过图来描述清楚了,这种方式的时间和为t2。发现t2比t1大很多。另外用nvprof测得每个kernel运行的时间也差异很大。请教下这可能得原因是什么呢,谢谢

上面提到的图指有向无环图,如下图所示,每个节点代表一个kernel函数运算得到的结果,每个节点之间的边描述了一个kernel函数。

另外测量时间用了两种方式,一种是cudaEventElapsedTime,包含了数据传输的时间。另一种是用nvprof测得的时间。

上面所述的两种执行方式所有相同kernel的输入参数均相同。


此外,对有向无环图的执行采用了两种编程模式,一种是数据都存在host端,每一个kernel运算都搬运数据过去,再把数据搬运回来。另外一种是数据存在device端,不需要中间结果的数据搬运,只需要在host端做一个存储管理单元(基本上不占时间)。但两种方式的结果都显示比单独执行的kernel时间长很多


页: [1]