여러 변수가 포함 된 사용자 정의 원자 함수를 구현하려면 어떻게해야합니까?

Question

나는 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 (어셈블리) 코드로 구현할 수 있습니까?

Answer 1

위의 두 번째 주석에서 언급했듯이 두 개의 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; 
} 

Answer 2

@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; 
}