gpt4 book ai didi

cuda - CUDA 扭曲调度是确定性的吗?

转载 作者:行者123 更新时间:2023-12-04 17:34:45 33 4
gpt4 key购买 nike

我想知道 CUDA 应用程序的扭曲调度顺序是否是确定性的。

具体来说,我想知道在同一设备上使用相同输入数据多次运行同一内核时,warp 执行的顺序是否会保持不变。如果没有,是否有任何东西可以强制对扭曲执行进行排序(例如在调试依赖于顺序的算法的情况下)?

最佳答案

CUDA 扭曲调度的精确行为没有定义。因此,您不能依赖它是确定性的。特别是,如果多个经线准备好在给定的发布槽中执行,则没有描述经线调度程序将选择哪个经线。

没有外部方法可以精确控制扭曲执行的顺序。

构建确定扭曲 ID 并强制扭曲以特定顺序执行的代码当然是可能的。像这样的东西:

#include <stdio.h>

#define N_WARPS 16
#define nTPB (32*N_WARPS)

__device__ volatile int my_next = 0;
__device__ int warp_order[N_WARPS];

__global__ void my_kernel(){

__shared__ volatile int warp_num;
unsigned my_warpid = (threadIdx.x & 0x0FE0U)>>5;
if (!threadIdx.x) warp_num = 0;
__syncthreads(); // don't use syncthreads() after this point
while (warp_num != my_warpid);
// warp specific code here
if ((threadIdx.x & 0x01F) == 0){
warp_order[my_next++] = my_warpid;
__threadfence();
warp_num++; // release next warp
} // could use syncthreads() after this point, if more code follows
}


int main(){

int h_warp_order[N_WARPS];
for (int i = 0; i < N_WARPS; i++) h_warp_order[i] = -1;
cudaMemcpyToSymbol(warp_order, h_warp_order, N_WARPS*sizeof(int));
my_kernel<<<1,nTPB>>>();
cudaDeviceSynchronize();
cudaMemcpyFromSymbol(h_warp_order, warp_order, N_WARPS*sizeof(int));
for (int i = 0; i < N_WARPS; i++) printf("index: %d, warp_id: %d\n", i, h_warp_order[i]);
return 0;
}

当然,一次只允许执行一个扭曲会非常低效。

一般来说,最好的可并行算法几乎没有或没有顺序依赖性。

关于cuda - CUDA 扭曲调度是确定性的吗?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/24977294/

33 4 0
Copyright 2021 - 2024 cfsdn All Rights Reserved 蜀ICP备2022000587号
广告合作:1813099741@qq.com 6ren.com