如何实现一个涉及多个变量的自定义原子函数？

我想CUDA来实现这个原子功能：如何实现一个涉及多个变量的自定义原子函数？

__device__ float lowest; // global var 
__device__ int lowIdx; // global var 
float realNum; // thread reg var 
int index;  // thread reg var 

if(realNum < lowest) { 
lowest= realNum; // the new lowest 
lowIdx= index; // update the 'low' index 
}

我不相信我可以与任何的原子功能做到这一点。我需要锁定几个全局内存loc的几条指令。我可以用PTXAS（汇编）代码来实现这个吗？

来源

2013-07-01 Doug

我不认为有一种方法（PTX或以其他方式）使用任何特定的GPU硬件一次自动更新多个位置。其他人可能有一个聪明的想法。通常情况下，我认为这种类型的问题可以使用“临界区”方法来解决，您可能想使用右上角的搜索框来搜索“cuda临界区”并查看其中一些问题中描述的内容。看来您可能希望以每个线程为基础运行此操作，并且每个线程临界区管理可能非常危险/困难。 –

实际上，对于这种有限的情况，你只有两个32位数量，你可能会创建一个自定义的原子函数，可能是围绕'atomicCAS'构建的，利用64位数量（巧妙地结合两个32位量），也许沿着[任意原子示例]给出的线（http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomic-functions）在文档中给出。 –

正如我在上面的第二条评论中所述，可以将两个32位数量组合成一个64位的原子管理数量，并以这种方式处理问题。然后，我们使用arbitrary atomic example作为粗略指南，以原子方式管理64位数量。显然你不能将这个想法扩展到两个32位数量。这里有一个例子：

#include <stdio.h> 
#define DSIZE 5000 
#define nTPB 256 

#define cudaCheckErrors(msg) \ 
    do { \ 
     cudaError_t __err = cudaGetLastError(); \ 
     if (__err != cudaSuccess) { \ 
      fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \ 
       msg, cudaGetErrorString(__err), \ 
       __FILE__, __LINE__); \ 
      fprintf(stderr, "*** FAILED - ABORTING\n"); \ 
      exit(1); \ 
     } \ 
    } while (0) 

typedef union { 
    float floats[2];     // floats[0] = lowest 
    int ints[2];      // ints[1] = lowIdx 
    unsigned long long int ulong; // for atomic update 
} my_atomics; 

__device__ my_atomics test; 

__device__ unsigned long long int my_atomicMin(unsigned long long int* address, float val1, int val2) 
{ 
    my_atomics loc, loctest; 
    loc.floats[0] = val1; 
    loc.ints[1] = val2; 
    loctest.ulong = *address; 
    while (loctest.floats[0] > val1) 
     loctest.ulong = atomicCAS(address, loctest.ulong, loc.ulong); 
    return loctest.ulong; 
} 


__global__ void min_test(const float* data) 
{ 

    int idx = (blockDim.x * blockIdx.x) + threadIdx.x; 
    if (idx < DSIZE) 
     my_atomicMin(&(test.ulong), data[idx],idx); 
} 

int main() { 

    float *d_data, *h_data; 
    my_atomics my_init; 
    my_init.floats[0] = 10.0f; 
    my_init.ints[1] = DSIZE; 

    h_data = (float *)malloc(DSIZE * sizeof(float)); 
    if (h_data == 0) {printf("malloc fail\n"); return 1;} 
    cudaMalloc((void **)&d_data, DSIZE * sizeof(float)); 
    cudaCheckErrors("cm1 fail"); 
    // create random floats between 0 and 1 
    for (int i = 0; i < DSIZE; i++) h_data[i] = rand()/(float)RAND_MAX; 
    cudaMemcpy(d_data, h_data, DSIZE*sizeof(float), cudaMemcpyHostToDevice); 
    cudaCheckErrors("cmcp1 fail"); 
    cudaMemcpyToSymbol(test, &(my_init.ulong), sizeof(unsigned long long int)); 
    cudaCheckErrors("cmcp2 fail"); 
    min_test<<<(DSIZE+nTPB-1)/nTPB, nTPB>>>(d_data); 
    cudaDeviceSynchronize(); 
    cudaCheckErrors("kernel fail"); 

    cudaMemcpyFromSymbol(&(my_init.ulong), test, sizeof(unsigned long long int)); 
    cudaCheckErrors("cmcp3 fail"); 

    printf("device min result = %f\n", my_init.floats[0]); 
    printf("device idx result = %d\n", my_init.ints[1]); 

    float host_val = 10.0f; 
    int host_idx = DSIZE; 
    for (int i=0; i<DSIZE; i++) 
    if (h_data[i] < host_val){ 
     host_val = h_data[i]; 
     host_idx = i; 
     } 

    printf("host min result = %f\n", host_val); 
    printf("host idx result = %d\n", host_idx); 
    return 0; 
}

来源

2013-07-01 21:27:46

一个绝妙的主意，非常感谢 – Doug

@Robert Crovella：很好的想法，但我认为函数应作如下修改一点点：

__device__ unsigned long long int my_atomicMin(unsigned long long int* address, float val1, int val2) 
{ 
    my_atomics loc, loctest, old; 
    loc.floats[0] = val1; 
    loc.ints[1] = val2; 
    loctest.ulong = *address; 
    old.ulong = loctest.ulong; 
    while (loctest.floats[0] > val1){ 
     old.ulong = loctest.ulong; 
     loctest.ulong = atomicCAS(address, loctest.ulong, loc.ulong); 
    } 
    return old.ulong; 
}

来源

2014-07-03 22:32:48

我不知道为什么。看来我们只是不同意函数的返回值。在您的情况下，返回值模式*与[文档中给出的示例]（http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html# atomic-functions），它返回* atomicCAS函数返回的最近值*（假设输入while循环）。你的品种不这样做。 –

如何实现一个涉及多个变量的自定义原子函数？

回答

相关问题