用户
 找回密码
 立即注册
发表于 2013-10-10 15:40:33
ice 发表于 2013-10-10 15:33
LZ您好:

我没说不能用for的,上面那个只是一个简单举例而已。

float c[16]={0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0};但是这样赋值为什么不行呢?
使用道具 举报 回复 支持 反对
发表于 2013-10-10 15:41:10
ice 发表于 2013-10-10 15:33
LZ您好:

我没说不能用for的,上面那个只是一个简单举例而已。

float c[16]={0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0};为什么这样赋值不行呢?分配到了局部存储器
使用道具 举报 回复 支持 反对
发表于 2013-10-10 15:52:20
wch8802 发表于 2013-10-10 15:41
float c[16]={0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0};为什么这样赋值不行呢?分配到了局部存储器 ...

请给出您的完整代码以便分析。
使用道具 举报 回复 支持 反对
发表于 2013-10-10 15:57:07
ice 发表于 2013-10-10 15:52
请给出您的完整代码以便分析。

__global__ void sgemmNN(const float* A,int lda,const float* B,int ldb,float* C,int ldc,int k)
{
        A+=blockIdx.x*64+threadIdx.x+threadIdx.y*16;
        B+=threadIdx.x+__mul24((__mul24(blockIdx.y,16)+threadIdx.y),ldb);
        C+=blockIdx.x*64+threadIdx.x+__mul24((threadIdx.y+__mul24(blockIdx.y,ldc)),16);
        int num=0;
        int id=blockIdx.y*64+threadIdx.y*16+threadIdx.x;

        __shared__ float bs[16][17];
        float c[16]={0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0};
        const float* Blast=B+k;
        do
        {
                //num++;
#pragma unroll
                for(int i=0;i<16;i+=4)
                        bs[threadIdx.x][threadIdx.y+i]=B[i*ldb];
                B+=16;
                __syncthreads();
#pragma unroll
                for(int i=0;i<16;i++,A+=lda)
                {
                        c[0]+=A[0]*bs[0];c[1]+=A[0]*bs[1];
                        c[2]+=A[0]*bs[2];c[3]+=A[0]*bs[3];
                        c[4]+=A[0]*bs[4];c[5]+=A[0]*bs[5];
                        c[6]+=A[0]*bs[6];c[7]+=A[0]*bs[7];
                        c[8]+=A[0]*bs[8];c[9]+=A[0]*bs[9];
                        c[10]+=A[0]*bs[10];c[11]+=A[0]*bs[11];
                        c[12]+=A[0]*bs[12];c[13]+=A[0]*bs[13];
                        c[14]+=A[0]*bs[14];c[15]+=A[0]*bs[15];

                }
                __syncthreads();
        }while(B<Blast);
        //pass[id]=num;
#pragma unroll
        for(int i=0;i<16;i++)
        {
                *(C+i*ldc)=c;
        }
}
使用道具 举报 回复 支持 反对
发表于 2013-10-10 16:07:13
ice 发表于 2013-10-10 15:52
请给出您的完整代码以便分析。

这种赋值我通过visual profiler 测出local memory有吞吐量
使用道具 举报 回复 支持 反对
发表于 2013-10-10 16:10:54
wch8802 发表于 2013-10-10 15:30
分别有两个关于local memory的选项,一个是local load(store) transactions,另一个是local load(store), ...

您可以参考profiler手册中的Matrics Reference章节的解释。


本帖子中包含更多资源

您需要 登录 才可以下载或查看,没有帐号?立即注册

x
使用道具 举报 回复 支持 反对
发表于 2013-10-10 16:31:18
ice 发表于 2013-10-10 16:10
您可以参考profiler手册中的Matrics Reference章节的解释。

经过我试验,数组都是分配到局部存储器(for循环也测出局部存储器吞吐量),当分别申请16个变量时,并以此赋值,局部存储器吞吐量就为0,我的问题是,如果数组很大,我申请很多变量就相当不方便,怎么解决这个问题呢,版主大人?
使用道具 举报 回复 支持 反对
发表于 2013-10-10 22:30:13
我来说两句,首先说,

楼主的“数组都将分配在local memory中"的说法是错误的。前文ICE版主也多次说明了。
这个很容易给出反例的,例如这种:
__global__ void wch(int *p)
{
     int sum[4] = {0, 0, 0, 0}; //int a = 0, b = 0, c = 0, d= 0;
     for (int i = 0; i < 10; i++)
     {
           sum[0] += p[i * 4 + 0]; //a += ....
           sum[1] += p[i * 4 + 1]; //b += ....
           sum[2] += p[i * 4 + 2]; //c += ....
           sum[3] += p[i * 4 + 3]; //d += ....
     }
    p[0] = sum[0] + sum[1] + sum[2] + sum[3]; //p[0] = a + b + c + d;
}
这个例子在-O2下无local traffic

其次,如果你分配100个float的数组,那是必然放不下的。这个请绝望(除非你在用3.5的卡,还是有可能的)。

最后,无法针对你的特定例子给出评价(因为此例子的代码一些下标被论坛吃掉了,请您重新看下您的发帖是否有问题),但如果在给定了足够的编译器限制的情况下,外加-O2编译,和你单独的16个变量无区别的。请重新尝试。

感谢来访。
使用道具 举报 回复 支持 反对
发表于 2013-10-10 22:33:38
横扫千军 发表于 2013-10-10 22:30
我来说两句,首先说,

楼主的“数组都将分配在local memory中"的说法是错误的。前文ICE版主也多次说明了。 ...

我发的上文有个文字错误,

最后一处的“编译器”应为“寄存器”。

特此更正。
使用道具 举报 回复 支持 反对
发表于 2013-10-22 16:37:15
横扫千军 发表于 2013-10-10 22:30
我来说两句,首先说,

楼主的“数组都将分配在local memory中"的说法是错误的。前文ICE版主也多次说明了。 ...

版主您好:如您上面提到的—O2编译是什么意思,我是直接进行编译的,怎么进行—O2编译
使用道具 举报 回复 支持 反对
您需要登录后才可以回帖 登录 | 立即注册