CUDA atomicCAS 使用方法

CUDA中提供的atomicCAS函数很有用,作为一个原子函数,能实现很多功能

首先,atomicCAS函数字样在VS中可能gcc不认识他,不要紧,nvcc认识CAS函数,所以虽然会报错,但写上去能跑

其次,atomicCAS函数不需要引用其他额外的头文件,但是这是个device函数,如果你要在host调用,请用一个global函数作为中介

其定义如下:

int atomicCAS(int* address, int compare, int val);

其中address指向的内存中变量值为old,这个函数返回当前old的值,然后如果old == compare,把address指向的内存中的值修改为val,如果old != compare,则保持address指向的值不变

即这个原子函数做了这样一件事:

int atomicCAS(int* address, int compare, int val)
{
    old = *address;
    if(old == compare)
        *address = val;
    else
        *address = old;
    return(old);
}

假设写了一个如下的代码:

int compare = 0;
int val = 1;
int old = 0;
int* address = &old
int output = atomicCAS(pt, B, C);

函数先比较address指针指向的变量old值(0)和compare的值(0),相同,然后将address指针指向的值改为val的值:1,返回本函数修改前的old值,即0

函数执行后old = 1,返回的output = 0

如果是这样:

int compare = 1;
int val = 1;
int old = 0;
int* address = &old
int output = atomicCAS(pt, B, C);

函数比较address指针指向变量old(0)和compare的值(1),发现不同,于是保持pt指针的值不变

函数执行后,old的值仍然为0,不变,返回的output = 0

利用atomicCAS可以对一个内存位置做一个线程锁,放置多线程同时访问一个地址。

如下,old是需要访问的数据,lock是和该内存位置配套的锁变量,true表示old变量此时没有被任何线程占用,可以被访问

bool lock = true;
int old = 100;

用atomicCAS 原子函数对该变量上锁

atomicCAS(&lock ,true ,false);

用atomicCAS原子函数解锁

atomicCAS(&lock ,false ,true);

为了让一个线程到了锁后,其他线程等待,用一个while函数持续争夺锁,直到锁被其他线程解开

while(!atomicCAS(&lock ,true ,false))
{
    continue;
}

以上函数就是不断询问,如果lock锁变量为false,即被占用,atomicCAS函数那么保持lock锁变量不变,返回lock原值false,同时while里条件判断为true,进行到下一轮询问

当锁被释放,lock锁变量变成true,atomicCAS函数将lock锁变量值修改成false,并且返回lock原值true,同时while里条件判断为false,跳出循环开始执行下面代码

完整代码:

while(!atomicCAS(&lock ,true ,false))
{
    continue;
}

old = old + 100;

sleep(1);

atomicCAS(&lock ,false ,true);

当两个线程同时执行时,假设线程1率先运行到while处,此时锁变量lock为true,atomicCAS将锁赋值为false,并且条件判断为!true,跳出循环,old = old +100,此时old值为100+100 =200

在线程1sleep期间,线程2也运行到了while处,此时锁变量已经被线程1修改为false了,atomicCAS不修改lock的值,条件判断为!false,继续循环。

直到线程1休眠完了,atomicCAS将lock值修改成了true,此时线程2还在循环while,突然发现lock值变成了true,于是while中的atomicCAS函数马上把lock值修改成false,条件判断为!true,跳出循环,执行old = old + 100,此时old值为200+100 = 300

需要注意的是,如果是block块内出现访问冲突,还是建议用共享内存修改,因为block块内的线程,每32个运行于同一个SM块,而SM块内线程是同步进行的。如此就会出现SM内一个线程得到了锁,其他线程都得不到锁,卡在那里,直接全SM锁死。

  • 3
    点赞
  • 3
    收藏
    觉得还不错? 一键收藏
  • 1
    评论
CUDA中的原子操作(atomic operation)是一种能够确保多个线程访问共享内存时的数据一致性的机制。原子操作可以保证某个特定内存位置的读取、修改或写入操作不会被其他线程中断,从而避免数据竞争问题。 在CUDA中,原子操作是通过特殊的函数来实现的,这些函数使用了硬件级别的原子指令。CUDA提供了不同类型的原子操作,如原子加法、原子减法、原子与、原子或等等。这些原子操作可以应用在不同的数据类型上,包括整型、浮点型和指针类型。 使用原子操作时需要注意以下几点: 1. 原子操作可能会引起较高的开销,因为它们需要保证多个线程之间的同步。因此,只有在必要时才应该使用原子操作。 2. 原子操作仅适用于共享内存,不能用于全局内存。 3. 原子操作不能保证全局的顺序一致性,只能保证特定内存位置的一致性。 下面是一个示例代码,展示了如何在CUDA使用原子操作进行加法: ```cpp __global__ void atomicAddKernel(int* data) { int tid = threadIdx.x + blockIdx.x * blockDim.x; atomicAdd(data, tid); } int main() { int data = 0; int* dev_data; cudaMalloc((void**)&dev_data, sizeof(int)); cudaMemcpy(dev_data, &data, sizeof(int), cudaMemcpyHostToDevice); atomicAddKernel<<<1, 256>>>(dev_data); cudaMemcpy(&data, dev_data, sizeof(int), cudaMemcpyDeviceToHost); cudaFree(dev_data); printf("Result: %d\n", data); // 输出结果应为 32640 return 0; } ```

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论 1
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值