最新帖子

页: [1] 2 3 ... 10
1
CUDA / dim3换线程组织形式导致无法调试的问题
« 最后发表 作者 轩令 昨天 10:12:31 pm »
kernel<<<32,32 >>>();
这种可以调试;
但改一下block里thread的组织结构后:
dim3 blockSize(32, 32);
kernel<<<32,blockSize >>> ();
这种无法调试,也就是核函数里的断点停不下来。

现在用的CUDA10.0,显卡是RTX2080ti,显卡驱动430.86。
2
CUDA / Re: 新手内存分配问题
« 最后发表 作者 jetlin1992 昨天 04:47:55 pm »
感谢补充!
但是楼主你的shared memory中的数组x_tmp[], 从填充了以后一次没用过啊!!!你如何能让我评判这种访存是否优劣?!

不能说什么了。楼主你自己去跑跑看看吧。

谢谢提醒,发现有两行写错了,在下面进行调用了

   Vi = x_tmp;//_根据前面两行获得的坐标查表获取数值
   Vj = x_tmp[j];//_根据前面两行获得的坐标查表获取数值
3
CUDA / Re: 新手内存分配问题
« 最后发表 作者 屠戮人神 昨天 04:39:09 pm »
抱歉,我对CUDA还不是很熟悉,发的内容可能不齐全,麻烦再帮我看看
具体代码如下。查表下标数据存在LineTable变量里面,根据总的进程ID Zindx 查到下标后,再去x_tmp里面找数据,最后算结果。
所有block共用同样的x_tmp,所以复制数据的时候没有用blockindx

程序代码: [选择]
#include "kernel.h"
#include <stdio.h>
 
//extern "C"
//Adds two arrays

#define NodeNum 3072
#define MemCpyNum 96 //3072节点,变量为3072*2 = 6144个元素,每个SM有64sp,每个线程复制96个数据
#define BlockD    512 //



__global__ void PQij_cal(const int* LineTable , float* P, float* Q, float* x)//LineTable线路两侧节点号表,x状态变量
{
__shared__ float x_tmp[6144];//总共6144个float

//__shared__ float line_para[];
//__shared__ int16_t ij;

//int i;
int Bindx = blockIdx.x; //block的编号
int Tindx = threadIdx.x;//每个block内的线程编号

//将数据从global memory 复制到shared memory
for (int i = 0; i < MemCpyNum; i++)//复制x数据
{
x_tmp[64 * i + Tindx] = x[64 * i + Tindx];//开64个thread
}
//
__syncthreads();


float Vi, Vj, AngD,SinA,CosA;
int Zindx, i, j,k;

for (k = 0; k < 8; k++)
{
Zindx = Bindx * BlockD + Tindx + 64 * k; //总id
i = LineTable[2 * Zindx]; //获得查表的下标
j = LineTable[2 * Zindx + 1]; //获得查表的下标

Vi = x[i];//_根据前面两行获得的坐标查表获取数值
Vj = x[j];//_根据前面两行获得的坐标查表获取数值

AngD = x[i + 3072] - x[j + 3072]; //_根据前面两行获得的坐标查表获取数值

P[Zindx] = Vi*Vj + AngD;
Q[Zindx] = Vi / Vj - AngD;//输出结果
}

};

感谢补充!
但是楼主你的shared memory中的数组x_tmp[], 从填充了以后一次没用过啊!!!你如何能让我评判这种访存是否优劣?!

不能说什么了。楼主你自己去跑跑看看吧。
4
CUDA / Re: 新手内存分配问题
« 最后发表 作者 jetlin1992 昨天 04:36:21 pm »
楼主你基本啥都没给出:
(1)是否“乱序”查询的时候”高效“,这个我不知道,因为你没有给出任何使用shared memory的代码,所以谈不上知道不知道”高效“使用了它与否。
(2)关于初始化shared memory中的内容的时候,你这种循环,没有给出64和MemCpyNum是什么。我们本着最好的情况说,如果64是你的block中的线程数量,而MemCpyNum是6K/64,则这个时候的载入是没有问题的。如果不是,则可能发生各种其他错误情况。
(3)直接使用这6K个float效果如何,这个我不能知道。在7.x之前的,5.x和6.x的卡,尽量要使用shared的,因为他们的L1 cache比较虚弱。但是你的7.5已经增强了L1 cache,这里可能会好一点,但具体如果,依然要看你的访存模型,如果是循环随机分散的,你依然还是需要上shared的,这个需要你实际的测试一下。
(4)如何使用L1/L2 cache,这个是自动的,不需要你操心,你正常的访存读取即可。如果L2/L1起作用,则自动的起作用了。这也是为何我们经常将Shared Memory叫做“手动cache”的原因,因为它不能自动,而L1/L2能。

此外,你发帖太不认真。你越认真,论坛就会越认真,这对你是好事。


关于L2cache的问题
在program guide里面有这么一段介绍

There is an L1 cache for each multiprocessor and an L2 cache shared by all
multiprocessors.
The L1 cache is used to cache accesses to local memory, including
temporary register spills. The L2 cache is used to cache accesses to local and global
memory. The cache behavior (e.g., whether reads are cached in both L1 and L2 or in
L2 only) can be partially configured on a per-access basis using modifiers to the load
or store instruction. Some devices of compute capability 3.5 and devices of compute
capability 3.7 allow opt-in to caching of global memory in both L1 and L2 via compiler
options.


这一块L2缓存,我看图灵有5MB那么多,我想问问是不是这个我们无法手动调用或者分配
5
CUDA / Re: 新手内存分配问题
« 最后发表 作者 jetlin1992 昨天 04:25:23 pm »
抱歉,我对CUDA还不是很熟悉,发的内容可能不齐全,麻烦再帮我看看
具体代码如下。查表下标数据存在LineTable变量里面,根据总的进程ID Zindx 查到下标后,再去x_tmp里面找数据,最后算结果。
所有block共用同样的x_tmp,所以复制数据的时候没有用blockindx

程序代码: [选择]
#include "kernel.h"
#include <stdio.h>
 
//extern "C"
//Adds two arrays

#define NodeNum 3072
#define MemCpyNum 96 //3072节点,变量为3072*2 = 6144个元素,每个SM有64sp,每个线程复制96个数据
#define BlockD    512 //



__global__ void PQij_cal(const int* LineTable , float* P, float* Q, float* x)//LineTable线路两侧节点号表,x状态变量
{
__shared__ float x_tmp[6144];//总共6144个float

//__shared__ float line_para[];
//__shared__ int16_t ij;

//int i;
int Bindx = blockIdx.x; //block的编号
int Tindx = threadIdx.x;//每个block内的线程编号

//将数据从global memory 复制到shared memory
for (int i = 0; i < MemCpyNum; i++)//复制x数据
{
x_tmp[64 * i + Tindx] = x[64 * i + Tindx];//开64个thread
}
//
__syncthreads();


float Vi, Vj, AngD,SinA,CosA;
int Zindx, i, j,k;

for (k = 0; k < 8; k++)
{
Zindx = Bindx * BlockD + Tindx + 64 * k; //总id
i = LineTable[2 * Zindx]; //获得查表的下标
j = LineTable[2 * Zindx + 1]; //获得查表的下标

Vi = x[i];//_根据前面两行获得的坐标查表获取数值
Vj = x[j];//_根据前面两行获得的坐标查表获取数值

AngD = x[i + 3072] - x[j + 3072]; //_根据前面两行获得的坐标查表获取数值

P[Zindx] = Vi*Vj + AngD;
Q[Zindx] = Vi / Vj - AngD;//输出结果
}

};
6
CUDA / Re: GPU SM中shared memory的调度
« 最后发表 作者 屠戮人神 昨天 03:54:35 pm »
(1)感谢您的回复,按照您的意思是说shared memory分配给多个block,且每个block分配的shared memory是独立且不同的,这些block中的shared memory加和就是我在代码中定义的总shared memory对吗?
(2)以及我在《CUDA 并行程序设计 编程指南》一书中看到如图的文字(图中的共享缓存指的是L2 Cache),如果按照书中所说,是否意味着L2 cache分配到所有SM的数据保持一致,每个SM的shared memory均相同。这么理解是否与(1)所述冲突?

(1)同一个kernel,每个block的shared memory都是独立的,他们是同样大小的多个独立副本。
(2)根据问题的不同,有的问题需要将一个大的区域拆分到每个block中的shared memory中,这个时候可以一定程度的认为“这些block用的shared memory总和加起来”是大问题所需要的总数据的结构;但是需要注意,有的问题就是需要将同样内容的shared memory,在每个block的shared memory,保持相同内容的只读副本,这个时候实际上每块shared memory本身就是完整的数据的结构的,并不存在需要“加和”才是的关系。例如本论坛刚刚有人提出需要一个Look-up table,放在shared memory中只读,此时完全可以每个人都是一个完整的副本的。

以上是对你的主要问题的讨论。

(3)关于你的第二部分,本论坛传统不为CSDN, 个人博客,知乎,其他非本公司出品的书籍(不含我司代理的NV或者华硕产品的),进行技术解答。请谅解。
7
CUDA / Re: 用cudaMalloc时一定要写(void**)吗?
« 最后发表 作者 屠戮人神 昨天 03:49:22 pm »
在用cudaMalloc函数分配设备内存时,一般类似这样写:
    float *a;
    int size=...
    cudaMalloc((void**)&a,size);
我想问,最后一行可以简化为如下形式吗?
    cudaMalloc(&a,size);
以上两种写法有无本质区别?

提前感谢解答。

在C里面,void*的是万能指针,这样是可以的。因为接受void *这种参数的,可以用其他任何指针给它。甚至还可以反过来,任何其他指针形式的,可以用void *这种万能的,不需要转换的给它赋值。这些可以直接用,无警告。但是C里面两个不同类型的指针直接给,是一个警告(但不一个错误)。

但是C++我就不确定了,我只知道C++里面不同类型的指针赋值将会是一个error,而不是一个warnings。而具体到能否将type*当成void*用,我就不清楚了。

欢迎补充C++中的情况,以便教育我。
8
CUDA / Re: GPU SM中shared memory的调度
« 最后发表 作者 appachi 昨天 03:38:58 pm »
(1)感谢您的回复,按照您的意思是说shared memory分配给多个block,且每个block分配的shared memory是独立且不同的,这些block中的shared memory加和就是我在代码中定义的总shared memory对吗?
(2)以及我在《CUDA 并行程序设计 编程指南》一书中看到如图的文字(图中的共享缓存指的是L2 Cache),如果按照书中所说,是否意味着L2 cache分配到所有SM的数据保持一致,每个SM的shared memory均相同。这么理解是否与(1)所述冲突?
9
CUDA / Re: 新手内存分配问题
« 最后发表 作者 屠戮人神 昨天 03:38:44 pm »
麻烦各位帮忙看看
需求:有一个常量数组约6000个 float ,需要在核函数中给各个线程乱序查询使用,想把这个数组放到shared memory里面,用下面的代码可以实现吗,会有内存冲突什么的吗。

还有一个问题,这个数组会给所有block共享,会不会放在L2 CACHE(RTX2080TI)里面更好,如果是的话,L2cache是怎么使用的?

__global__ void PQij_cal(const int* LineTable , float* LineV, float* LineAng, float* x)//LineTable线路两侧节点号表,x状态变量
{
   __shared__ float x_tmp[6144];//总共6144个float
   __shared__ float line_para[];
   __shared__ int16_t ij;

   //int i;
   int Bindx = blockIdx.x; //block的编号
   int Tindx = threadIdx.x;//每个block内的线程编号
   //将数据从global memory 复制到shared memory
   for (int i = 0; i < MemCpyNum; i++)//复制x数据
   {
      xtmp[64 * i + Tindx] = x[64 * i + Tindx];//最后一个下标为
   }
   //
   __syncthreads();

   //计算PQ

楼主你基本啥都没给出:
(1)是否“乱序”查询的时候”高效“,这个我不知道,因为你没有给出任何使用shared memory的代码,所以谈不上知道不知道”高效“使用了它与否。
(2)关于初始化shared memory中的内容的时候,你这种循环,没有给出64和MemCpyNum是什么。我们本着最好的情况说,如果64是你的block中的线程数量,而MemCpyNum是6K/64,则这个时候的载入是没有问题的。如果不是,则可能发生各种其他错误情况。
(3)直接使用这6K个float效果如何,这个我不能知道。在7.x之前的,5.x和6.x的卡,尽量要使用shared的,因为他们的L1 cache比较虚弱。但是你的7.5已经增强了L1 cache,这里可能会好一点,但具体如果,依然要看你的访存模型,如果是循环随机分散的,你依然还是需要上shared的,这个需要你实际的测试一下。
(4)如何使用L1/L2 cache,这个是自动的,不需要你操心,你正常的访存读取即可。如果L2/L1起作用,则自动的起作用了。这也是为何我们经常将Shared Memory叫做“手动cache”的原因,因为它不能自动,而L1/L2能。

此外,你发帖太不认真。你越认真,论坛就会越认真,这对你是好事。
10
CUDA / 新手内存分配问题
« 最后发表 作者 jetlin1992 昨天 02:54:29 pm »
麻烦各位帮忙看看
需求:有一个常量数组约6000个 float ,需要在核函数中给各个线程乱序查询使用,想把这个数组放到shared memory里面,用下面的代码可以实现吗,会有内存冲突什么的吗。

还有一个问题,这个数组会给所有block共享,会不会放在L2 CACHE(RTX2080TI)里面更好,如果是的话,L2cache是怎么使用的?

__global__ void PQij_cal(const int* LineTable , float* LineV, float* LineAng, float* x)//LineTable线路两侧节点号表,x状态变量
{
   __shared__ float x_tmp[6144];//总共6144个float
   __shared__ float line_para[];
   __shared__ int16_t ij;

   //int i;
   int Bindx = blockIdx.x; //block的编号
   int Tindx = threadIdx.x;//每个block内的线程编号
   //将数据从global memory 复制到shared memory
   for (int i = 0; i < MemCpyNum; i++)//复制x数据
   {
      xtmp[64 * i + Tindx] = x[64 * i + Tindx];//最后一个下标为
   }
   //
   __syncthreads();

   //计算PQ
   
页: [1] 2 3 ... 10