最新帖子

页: 1 [2] 3 4 ... 10
11
CUDA / SM、block、warp
« 最后发表 作者 jinyer 四月 01, 2022, 03:56:39 pm »
一个Multi-Processor上面可以同时驻留多个warps,例如一个计算能力8.6的卡,它的1个SM里面,最多可能驻留1536个线程,也就是大约48个warps,在同时等待被调度执行。并且同时最多能上16个block。
请问一个SM是同时在执行16个block吗?
若是的话,SM一个时刻只能执行每个block里边的一个warp吗?即并行线程数为 82*16*32吗?【目前我的理解是这样,可3090卡的SP个数只有10496是82*16*32的四分之一,这又是怎么做到的呢】
12
CUDA / Re: race condition结果一致
« 最后发表 作者 屠戮人神 三月 30, 2022, 04:02:24 pm »
我将kernel循坏执行100次,指定计算架构为X=7, Y=0。
下面的代码输出s的值全为2,所以该warp内的线程是先共同请求s, 放于register file, 然后同时将数据拿到计算单元吗。?
但是我在您的另一篇答疑中看到您表示可能存在s[0]+=s[16] 和s[16]+=s[32]次序的问题。
是写回速度相对较慢导致难以出现s[0]加已经+=s[32]的s[16]的值吗?
程序代码: [选择]
__global__ void kernel(){
    int tid=threadIdx.x;
    extern __shared__ int s[];
    s[tid]=1;
    s[32]=1;
    __syncwarp();
    s[tid] += s[tid+1];
    printf("%d %d\n",s[tid],tid);
}
int main(void){
    for(int i=0;i<100;i++){
        kernel<<<1,32,132>>>();
        cudaDeviceSynchronize();
    }
}
希望您能帮我解答。如果您需要的话,可以奉上输出的内容。

在下面的代码中,volatile和被注释掉的__syncwarp()效果相同,都为正确结果。编译时指定的架构为X=7,Y=0。
volatile 自带同步功能吗?还是70架构仍可以 lock_steped?
程序代码: [选择]
#include<stdio.h>
__global__ void kernel(){
    int tid=threadIdx.x;
    extern __shared__ volatile int s[];
    s[tid]=1;
    s[tid+8]=1;
    __syncwarp();
    s[tid] += s[tid+1];
    // __syncwarp();
    s[tid] += s[tid+4];
    // __syncwarp();
    s[tid] += s[tid+2];
    printf("%d %d\n",s[tid],tid);
}
int main(void){
    for(int i=0;i<100;i++){
        kernel<<<1,32,160>>>();
        cudaDeviceSynchronize();
    }
}

我感觉你的思维逻辑出现问题:
很多东西有多次证明,并不能说它成立(因为你将来可能遭遇到任意一次不成立的,就可以说明它不对了)。

无更多可用解释。我无法从你的演示代码中得到“Volta+依然是locked step执行warp”的任何结论。
13
CUDA / race condition结果一致
« 最后发表 作者 Dragon_ 三月 25, 2022, 03:37:08 pm »

我将kernel循坏执行100次,指定计算架构为X=7, Y=0。
下面的代码输出s的值全为2,所以该warp内的线程是先共同请求s, 放于register file, 然后同时将数据拿到计算单元吗。?
但是我在您的另一篇答疑中看到您表示可能存在s[0]+=s[16] 和s[16]+=s[32]次序的问题。
是写回速度相对较慢导致难以出现s[0]加已经+=s[32]的s[16]的值吗?
程序代码: [选择]
__global__ void kernel(){
    int tid=threadIdx.x;
    extern __shared__ int s[];
    s[tid]=1;
    s[32]=1;
    __syncwarp();
    s[tid] += s[tid+1];
    printf("%d %d\n",s[tid],tid);
}
int main(void){
    for(int i=0;i<100;i++){
        kernel<<<1,32,132>>>();
        cudaDeviceSynchronize();
    }
}
希望您能帮我解答。如果您需要的话,可以奉上输出的内容。

在下面的代码中,volatile和被注释掉的__syncwarp()效果相同,都为正确结果。编译时指定的架构为X=7,Y=0。
volatile 自带同步功能吗?还是70架构仍可以 lock_steped?
程序代码: [选择]
#include<stdio.h>
__global__ void kernel(){
    int tid=threadIdx.x;
    extern __shared__ volatile int s[];
    s[tid]=1;
    s[tid+8]=1;
    __syncwarp();
    s[tid] += s[tid+1];
    // __syncwarp();
    s[tid] += s[tid+4];
    // __syncwarp();
    s[tid] += s[tid+2];
    printf("%d %d\n",s[tid],tid);
}
int main(void){
    for(int i=0;i<100;i++){
        kernel<<<1,32,160>>>();
        cudaDeviceSynchronize();
    }
}
14
CUDA / Re: 请教:现在较新的GPU规定一个寄存器内存大小了吗?
« 最后发表 作者 cudaxiaox 三月 23, 2022, 03:19:00 pm »
我猜想你的问题的意思是每个线程可以最多使用多少个4B的常规寄存器?根据手册,你指出的两款卡都是可以使用255个的。

但是你的提问方式“寄存器上的变量占用多大内存?”的确很令人费解。

谢谢,我明白了!
15
CUDA / Re: 请教:现在较新的GPU规定一个寄存器内存大小了吗?
« 最后发表 作者 屠戮人神 三月 22, 2022, 09:27:00 pm »
请问像TITAN V、RTX3090这些显卡有没有规定CUDA编程时一个线程内可分配的存储在寄存器上的变量最多占多大内存呢?

我猜想你的问题的意思是每个线程可以最多使用多少个4B的常规寄存器?根据手册,你指出的两款卡都是可以使用255个的。

但是你的提问方式“寄存器上的变量占用多大内存?”的确很令人费解。
16
CUDA / Re: 同一个CUDA流操作
« 最后发表 作者 屠戮人神 三月 22, 2022, 09:19:59 pm »
看到大家说:同一个CUDA流操作是严格顺序的。
但是我跑代码时却出现了不一样的结果
程序代码: [选择]
noisePdf1<<<gridNoiPdf1, blockNoiPdf1 >>>(noiPdf1, d_rx, sigma);
printf("After---------------noisePdf1\n");
convPy1<<<gridConvPy, blockConvPy >>>(convResult, u0, u1);
printf("After------------------convPy\n");
然后分别在两个核函数中有一个printf语句。
在默认流中先后调用这两个核函数,结果输出语句如下:
After---------------noisePdf1
After------------------convPy
run----noisePdf1
running------convPy
请问这是表示如果第二个核函数想用第一个核函数的结果,这两个核函数调用中间必须加cudaDeviceSynchronize();吗?

你在偷换概念, 每个流中的任务(kernel计算任务,或者数据传输任务)的确是按照顺序进行的。但是这不代表,夹杂在中间的其他CPU上的操作,也要按照GPU的流中的顺序。

实际上你的这个例子是,每次kernel启动后,CPU不必等待kernel执行完毕,就可以继续往下执行了(kernel的启动是异步的)。所以你看到这样的是正常的。

只有计算和传输任务才会在同1个流中(你的例子中是默认流,我建议你使用非默认流避免你遭遇其他情况的问题)按照你发布命令的顺序执行的,例如你的第二问中的,连续2个kernel,依次顺序执行,这点是没有问题的。
17
CUDA / 请教:现在较新的GPU规定一个寄存器内存大小了吗?
« 最后发表 作者 cudaxiaox 三月 21, 2022, 08:53:38 am »
请问像TITAN V、RTX3090这些显卡有没有规定CUDA编程时一个线程内可分配的存储在寄存器上的变量最多占多大内存呢?
18
CUDA / Re: 同一个CUDA流操作
« 最后发表 作者 jinyer 三月 16, 2022, 02:58:50 pm »
是异步调度,但实际执行的时候还是要等第一个kernel执行完再会执行第二个kernel?或者是说 并行资源足够的情况下,两个核函数会一起执行呢
19
CUDA / 同一个CUDA流操作
« 最后发表 作者 jinyer 三月 16, 2022, 11:37:08 am »
看到大家说:同一个CUDA流操作是严格顺序的。
但是我跑代码时却出现了不一样的结果
程序代码: [选择]
noisePdf1<<<gridNoiPdf1, blockNoiPdf1 >>>(noiPdf1, d_rx, sigma);
printf("After---------------noisePdf1\n");
convPy1<<<gridConvPy, blockConvPy >>>(convResult, u0, u1);
printf("After------------------convPy\n");
然后分别在两个核函数中有一个printf语句。
在默认流中先后调用这两个核函数,结果输出语句如下:
After---------------noisePdf1
After------------------convPy
run----noisePdf1
running------convPy
请问这是表示如果第二个核函数想用第一个核函数的结果,这两个核函数调用中间必须加cudaDeviceSynchronize();吗?
20
有可能是这个原因。但是去掉了volatile后永远不会检测到正确的flag了。

实际上有没有volatile是两种指令,
一种是:
LDG.STRONG.SYS
另外一种:LDG
所以可能L2也会缓冲系统内存(因为后一种多次读取都检测不到,可能是kernel运行期间将旧值缓冲了)。
嗯嗯从结果来看是会缓冲的
页: 1 [2] 3 4 ... 10