OpenCL中能否用布尔缓冲区在核执行时修改值终止循环？执行挂起求助

我想在OpenCL环境下做个实验，验证能不能通过主机代码修改缓冲区，在内核执行过程中改变其运行状态。尝试用这种方式终止内核里的while循环时，程序出现了挂起现象。

内核代码 #

void my_kernel(
                __global bool *in,
                __global int *out)
{
        int i = get_global_id(0);
        while(1) {
                if(1 == *in) {
                        printf("while loop is finished");
                        break;
                }
        }
        printf("out[0] = %d\n", out[0]);
}

主机修改缓冲区的代码 #

input[0] = 1; 
err = clEnqueueWriteBuffer(commands, input_buffer,
                CL_TRUE, 0, sizeof(int), (void*)input,
                0, NULL,NULL);

问题原因分析 #

1. 编译器优化导致读取缓存值：内核中__global bool *in指向的变量没有被内核写入，编译器会默认把它的值缓存到寄存器里。这就导致循环里每次读取的都是初始的缓存值（0），永远检测不到主机修改后的1，循环无法终止，程序挂起。
2. 内存类型不匹配：内核里in是bool类型（1字节），但主机代码中用sizeof(int)（4字节）写入数据，会造成内存越界，可能破坏其他内存区域，引发不可预料的问题。

解决方案 #

1. 添加volatile限定符禁止缓存：给in指针加上volatile，强制编译器每次都从全局内存读取最新值，而不是用寄存器缓存：

void my_kernel(
                __global volatile bool *in,
                __global int *out)
{
        int i = get_global_id(0);
        while(1) {
                if(1 == *in) {
                        printf("while loop is finished");
                        break;
                }
        }
        printf("out[0] = %d\n", out[0]);
}

2. 修正内存操作的类型匹配：主机代码中把input的类型改为bool，写入长度改为sizeof(bool)，避免内存越界：

bool input[1] = {0};
// 初始化缓冲区等前置操作
input[0] = 1; 
err = clEnqueueWriteBuffer(commands, input_buffer,
                CL_TRUE, 0, sizeof(bool), (void*)input,
                0, NULL,NULL);

内容的提问来源于stack exchange，提问作者Tigran84