가우스 제거 병렬 처리

Question

가우스 제거를 위해 CUDA에서 단일 스레드 프로그램을 성공적으로 구현했으며 병렬 처리를 달성하고자합니다. 지금까지 병렬 코드를 보이는 같은 :

__global__ void ParallelGaussian(double* A) 
{ 
    int index = threadIdx.x; 
    int stride = blockDim.x; 

    if (index < ROWS) //Skip additional threads 
    { 
     for (unsigned int r = index; r < ROWS; r += stride) 
     { 
      //Forward elimination to reduce to row echelon form 
      for (unsigned int k = r + 1; k < ROWS; ++k) 
      { 
       double c = -A[(ROWS + 1) * k + r]/A[(ROWS + 1) * r + r]; 
       for (unsigned int j = r; j < ROWS + 1; ++j) 
       { 
        if (r == j) 
         A[(ROWS + 1) * k + j] = 0.0; 
        else 
         A[(ROWS + 1) * k + j] += c * A[(ROWS + 1) * r + j]; 
       } 
      } 
     } 
    } 
} 

우리는 낮은 삼각 행렬에 1 차원 배열 (매트릭스)를 변환 할 GPU의 코드를 볼 수 있습니다 후 CPU에 내가 다시 계속 하겠지만 대체하여 최종 결과를 얻습니다. 완전히 필요하지는 않지만 실제로 알고리즘의 수치 안정성을 향상시키기 때문에이 접근법에서는 피벗이 수행되지 않습니다.

단일 스레드 블록 작품으로 커널을 시작하고 행에 쉴론 형태로 행렬을 변환 : 나는 스레드의 수를 증가하고자하는 경우

ParallelGaussian << < 1, 32 >> >(dev_a); 

처럼,

ParallelGaussian << < 1, 1 >> >(dev_a); 

그러나

낮은 삼각 행렬을 생성하지 못합니다. 이제 __syncthreads() 호출을 코드에 추가하여 블록의 스레드를 동기화해도 상황이 개선되지 않고 왜 그럴 수 없습니까?

Answer 1

내부 루프를 고려하십시오. 모든 스레드는 A에 액세스하며 k과 j이 행의 끝에 이르기까지 r에서 실행되기 때문에 복수 스레드가 동일한 A[(ROWS + 1) * k + j] 값을 수정할 가능성이 있습니다.

다른 스레드가 해당 값을 업데이트하는 동안 A[(ROWS + 1) * r + j]에 액세스하는 일부 스레드가있을 수도 있습니다.

하나의 가능한 솔루션은 각 스레드를 개별 결과 배열에 누적시킨 다음 끝에 결합하는 것입니다. 이것은 메모리 집약적입니다.

또 하나의 스레드가 특정 값에 쓰고 새 행렬에 값을 저장하여 다른 스레드가 필요로하는 값을 변경하지 않도록 구조를 변경하는 것입니다.