•  12-11-2019
  •  | 
  •  

문제

다양한 블록에서 최종 합계의 원자 계산이있는 이중 정밀 배열을 위해 클래식 Dot-Product 커널을 구현하려고합니다.Programming Guide의 116 페이지에 명시된대로 AtomicAdadd를 사용했습니다. 필자는 잘못된 것을하고 있습니다. 모든 블록의 스레드를 가로 지르는 부분 합계는 올바르게 계산되지만 Atowners는 AtoWords가 제대로 작동하지 않는 것 같습니다.동일한 데이터로 커널을 실행할 때마다 다른 결과가 나타납니다.누군가가 실수를 발견하거나 대체 솔루션을 제공 할 수 있다면 감사 할 것입니다! 여기 내 커널이 있습니다 : 

__global__ void cuda_dot_kernel(int *n,double *a, double *b, double *dot_res)
{
    __shared__ double cache[threadsPerBlock]; //thread shared memory
    int global_tid=threadIdx.x + blockIdx.x * blockDim.x;
    int i=0,cacheIndex=0;
    double temp = 0;
    cacheIndex = threadIdx.x;
    while (global_tid < (*n)) {
        temp += a[global_tid] * b[global_tid];
        global_tid += blockDim.x * gridDim.x;
    }
    cache[cacheIndex] = temp;
    __syncthreads();
    for (i=blockDim.x/2; i>0; i>>=1) {
        if (threadIdx.x < i) {
            cache[threadIdx.x] += cache[threadIdx.x + i];
        }
        __syncthreads();
    }
    __syncthreads();
    if (cacheIndex==0) {
        *dot_res=cuda_atomicAdd(dot_res,cache[0]);
    }
}
.

여기에서 내 장치 기능은 AtomicAdd : 

__device__ double cuda_atomicAdd(double *address, double val)
{
    double assumed,old=*address;
    do {
        assumed=old;
        old= __longlong_as_double(atomicCAS((unsigned long long int*)address,
                    __double_as_longlong(assumed),
                    __double_as_longlong(val+assumed)));
    }while (assumed!=old);

    return old;
}
.

도움이 되었습니까?

해결책

cuda_atomicAdd 기능을 잘못 사용하고 있습니다.커널 의이 섹션 : 

if (cacheIndex==0) {
    *dot_res=cuda_atomicAdd(dot_res,cache[0]);
}
.

는 범인입니다.여기에서는 dot_res에 원자 적으로 추가합니다.그런 다음 비 원자로 세트 dot_res를 결과로 설정합니다.이 함수의 리턴 결과는 원자 적으로 업데이트되는 위치의 이전 값 이며 "정보"또는 호출자의 로컬 사용에만 제공됩니다.원자 적으로 업데이트 된 것에 할당하지 않아서 처음에는 원자력 메모리 액세스를 사용하는 목적을 완전히 패배시킵니다.대신 이런 일을하십시오 : 

if (cacheIndex==0) {
    double result=cuda_atomicAdd(dot_res,cache[0]);
}
.

다른 팁

Ad Hoc CUDA 코드를 사용하여 감소 할 수있는 오른쪽은 까다로울 수 있으므로 CUDA 툴킷에 포함 된 추력 알고리즘을 사용하는 대체 솔루션입니다. 

#include <thrust/inner_product.h>
#include <thrust/device_ptr.h>

double do_dot_product(int n, double *a, double *b)
{
  // wrap raw pointers to device memory with device_ptr
  thrust::device_ptr<double> d_a(a), d_b(b);

  // inner_product implements a mathematical dot product
  return thrust::inner_product(d_a, d_a + n, d_b, 0.0);
}
.

코드가 해당 깊이를 확인하지 않았지만 여기에는 조언이 있습니다.
복잡한 문제가 발생할 경우 GPU를 효율적으로 프로그래밍할지 모르겠다면 Complex Protuct를 사용하면 GPU만을 사용하는 경우에만 추진을 사용하여 조언을 사용합니다.

  1. 도트 제품을 요약하도록 새로운 병렬 감소 커널을 시작하십시오.
    데이터가 이미 장치에 있기 때문에 새 커널을 시작하는 성능 저하가 표시되지 않습니다.

  2. 커널은 최신 GPU에서 최대한의 가능한 블록 수를 가로 질러 확장하지 않는 것 같습니다. 커널이 수백만 가지 값의 도트 산물을 계산할 수 있고 직렬화 된 원자 작업으로 인해 성능이 크게 감소 할 수 있습니다.

  3. 초보자 실수 : 입력 데이터와 공유 메모리 액세스 범위가 을 선택합니까? 또는 입력 데이터가 항상 블록 크기의 배수 인 것입니까? 그렇지 않으면 쓰레기를 읽습니다. 내 잘못된 결과의 대부분은이 잘못 때문입니다.

  4. 병렬 감소를 최적화합니다. 내 논문 또는 최적화 Mark Harris

    테스트되지 않은, 방금 메모장에서 썼습니다. 

    /*
 * @param inCount_s unsigned long long int Length of both input arrays
 * @param inValues1_g double* First value array
 * @param inValues2_g double* Second value array
 * @param outDots_g double* Output dots of each block, length equals the number of blocks
 */
__global__ void dotProduct(const unsigned long long int inCount_s,
    const double* inValuesA_g,
    const double* inValuesB_g,
    double* outDots_g)
{
    //get unique block index in a possible 3D Grid
    const unsigned long long int blockId = blockIdx.x //1D
            + blockIdx.y * gridDim.x //2D
            + gridDim.x * gridDim.y * blockIdx.z; //3D


    //block dimension uses only x-coordinate
    const unsigned long long int tId = blockId * blockDim.x + threadIdx.x;

    /*
     * shared value pair products array, where BLOCK_SIZE power of 2
     *
     * To improve performance increase its size by multiple of BLOCK_SIZE, so that each threads loads more then 1 element!
     * (outDots_g length decreases by same factor, and you need to range check and initialize memory)
     * -> see harris gpu optimisations / parallel reduction slides for more informations.
     */
    __shared__ double dots_s[BLOCK_SIZE];


    /*
     * initialize shared memory array and calculate dot product of two values, 
     * shared memory always needs to be initialized, its never 0 by default, else garbage is read later!
     */
    if(tId < inCount_s)
        dots_s[threadIdx.x] = inValuesA_g[tId] * inValuesB_g[tId];
    else
        dots_s[threadIdx.x] = 0;
    __syncthreads();

    //do parallel reduction on shared memory array to sum up values
    reductionAdd(dots_s, dots_s[0]) //see my thesis link

    //output value
    if(threadIdx.x == 0)
        outDots_g[0] = dots_s[0];

    //start new parallel reduction kernel to sum up outDots_g!
}
    .

    편집 : 불필요한 점을 제거했습니다.

라이센스 : CC-BY-SA ~와 함께 속성
제휴하지 않습니다 StackOverflow
scroll top