sách gpt4 ai đã đi

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

In lại 作者:行者123 更新时间:2023-12-04 17:34:45 33 4
mua khóa gpt4 Nike

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

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

1 Câu trả lời

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

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

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

#include 

#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]);
trả về 0;
}

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

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

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

33 4 0
行者123
Hồ sơ cá nhân

Tôi là một lập trình viên xuất sắc, rất giỏi!

Nhận phiếu giảm giá Didi Taxi miễn phí
Mã giảm giá Didi Taxi
Giấy chứng nhận ICP Bắc Kinh số 000000
Hợp tác quảng cáo: 1813099741@qq.com 6ren.com