문제
다양한 블록에서 최종 합계의 원자 계산이있는 이중 정밀 배열을 위해 클래식 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만을 사용하는 경우에만 추진을 사용하여 조언을 사용합니다.
-
도트 제품을 요약하도록 새로운 병렬 감소 커널을 시작하십시오.
데이터가 이미 장치에 있기 때문에 새 커널을 시작하는 성능 저하가 표시되지 않습니다. -
커널은 최신 GPU에서 최대한의 가능한 블록 수를 가로 질러 확장하지 않는 것 같습니다. 커널이 수백만 가지 값의 도트 산물을 계산할 수 있고 직렬화 된 원자 작업으로 인해 성능이 크게 감소 할 수 있습니다.
-
초보자 실수 : 입력 데이터와 공유 메모리 액세스 범위가 을 선택합니까? 또는 입력 데이터가 항상 블록 크기의 배수 인 것입니까? 그렇지 않으면 쓰레기를 읽습니다. 내 잘못된 결과의 대부분은이 잘못 때문입니다.
-
병렬 감소를 최적화합니다. 내 논문 또는 최적화 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! }
편집 : 불필요한 점을 제거했습니다.