race condition结果一致

  • 1 replies
  • 1080 views
race condition结果一致
« 于: 三月 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();
    }
}

Re: race condition结果一致
« 回复 #1 于: 三月 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”的任何结论。