CUDA : 모든 벡터 요소의 계단식 합계

Question

GPU와 CPU에서 float 값의 큰 벡터에 대해 계단식 덧셈 기능을 구현했습니다. 이는 단순히이 벡터 껍질의 모든 요소가 하나의 결과로 합산된다는 것을 의미합니다. CPU 알고리즘은 매우 사소하고 잘 작동하지만, GPU 알고리즘은 항상 원하는 결과에서 35200입니다.

알고리즘과 CPU 비교를위한 최소 작업 코드는 아래와 같습니다.

출력은 항상 이것이다 :

CPU Time: 22.760059 ms, bandwidth: 3.514929 GB/s 

GPU Time (improved): 12.077088 ms, bandwidth: 6.624114 GB/s 
- CPU result does not match GPU result in improved atomic add. 
    CPU: 10000000.000000, GPU: 10035200.000000, diff:-35200.000000 

나는 CUDA-memcheck 그것을 확인하지만 오류는 실행에 발생합니다. 나는 많은 다른 것들을 시도했지만 그 중 아무 것도 시도하지 못했습니다. 모든 float를 int로 변경했기 때문에 float 데이터 유형의 부정확성으로 인한 것이 아닌데도 똑같은 결과가 나타납니다. 난 당신이 커널 호출에서 중복 인덱스 생각

#include "cuda_runtime.h" 
#include "device_launch_parameters.h" 

#include <chrono> 
#include <time.h> 
#include <stdio.h> 
#include <stdlib.h> 

void reductionWithCudaImproved(float *result, const float *input); 
__global__ void reductionKernelImproved(float *result, const float *input); 
void reductionCPU(float *result, const float *input); 

#define SIZE 10000000 

#define TILE 32 

#define ILP 8 
#define BLOCK_X_IMPR (TILE/ILP) 
#define BLOCK_Y_IMPR 32 
#define BLOCK_COUNT_X_IMPR 100 

int main() 
{ 
    int i; 
    float *input; 
    float resultCPU, resultGPU; 
    double cpuTime, cpuBandwidth; 

    input = (float*)malloc(SIZE * sizeof(float)); 
    resultCPU = 0.0; 
    resultGPU = 0.0; 

    srand((int)time(NULL)); 

    auto start = std::chrono::high_resolution_clock::now(); 
    auto end = std::chrono::high_resolution_clock::now(); 

    for (i = 0; i < SIZE; i++) 
     input[i] = 1.0; 

    start = std::chrono::high_resolution_clock::now(); 
    reductionCPU(&resultCPU, input); 
    end = std::chrono::high_resolution_clock::now(); 

    std::chrono::duration<double> diff = end - start; 
    cpuTime = (diff.count() * 1000); 
    cpuBandwidth = (sizeof(float) * SIZE * 2)/(cpuTime * 1000000); 
    printf("CPU Time: %f ms, bandwidth: %f GB/s\n\n", cpuTime, cpuBandwidth); 

    reductionWithCudaImproved(&resultGPU, input); 

    if (resultCPU != resultGPU) 
     printf("- CPU result does not match GPU result in improved atomic add. CPU: %f, GPU: %f, diff:%f\n\n", resultCPU, resultGPU, (resultCPU - resultGPU)); 
    else 
     printf("+ CPU result matches GPU result in improved atomic add. CPU: %f, GPU: %f\n\n", resultCPU, resultGPU); 

    return 0; 
} 

void reductionCPU(float *result, const float *input) 
{ 
    for (int i = 0; i < SIZE; i++) 
     *result += input[i]; 
} 

__global__ void reductionKernelImproved(float *result, const float *input) 
{ 
    int i; 
    int col = (blockDim.x * blockIdx.x + threadIdx.x) * ILP; 
    int row = blockDim.y * blockIdx.y + threadIdx.y; 
    int index = row * blockDim.x * BLOCK_COUNT_X_IMPR + col; 
    __shared__ float interResult; 

    if (threadIdx.x == 0 && threadIdx.y == 0) 
     interResult = 0.0; 

    __syncthreads(); 

#pragma unroll ILP 
    for (i = 0; i < ILP; i++) 
    { 
     if (index < SIZE) 
     { 
      atomicAdd(&interResult, input[index]); 
      index++; 
     } 
    } 

    __syncthreads(); 

    if (threadIdx.x == 0 && threadIdx.y == 0) 
     atomicAdd(result, interResult); 
} 

void reductionWithCudaImproved(float *result, const float *input) 
{ 
    dim3 dim_grid, dim_block; 

    float *dev_input = 0; 
    float *dev_result = 0; 
    cudaEvent_t start, stop; 
    float elapsed = 0; 
    double gpuBandwidth; 

    dim_block.x = BLOCK_X_IMPR; 
    dim_block.y = BLOCK_Y_IMPR; 
    dim_block.z = 1; 

    dim_grid.x = BLOCK_COUNT_X_IMPR; 
    dim_grid.y = (int)ceil((float)SIZE/(float)(TILE * dim_block.y* BLOCK_COUNT_X_IMPR)); 
    dim_grid.z = 1; 

    cudaSetDevice(0); 

    cudaMalloc((void**)&dev_input, SIZE * sizeof(float)); 
    cudaMalloc((void**)&dev_result, sizeof(float)); 
    cudaMemcpy(dev_input, input, SIZE * sizeof(float), cudaMemcpyHostToDevice); 
    cudaMemcpy(dev_result, result, sizeof(float), cudaMemcpyHostToDevice); 

    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 

    cudaEventRecord(start); 
    reductionKernelImproved << <dim_grid, dim_block >> >(dev_result, dev_input); 

    cudaEventRecord(stop); 
    cudaEventSynchronize(stop); 

    cudaEventElapsedTime(&elapsed, start, stop); 

    gpuBandwidth = (sizeof(float) * SIZE * 2)/(elapsed * 1000000); 
    printf("GPU Time (improved): %f ms, bandwidth: %f GB/s\n", elapsed, gpuBandwidth); 

    cudaDeviceSynchronize(); 

    cudaMemcpy(result, dev_result, sizeof(float), cudaMemcpyDeviceToHost); 

    cudaFree(dev_input); 
    cudaFree(dev_result); 

    return; 
} 

Answer 1

:

이

내 코드 내가 잘못 아니에요 경우

int col = (blockDim.x * blockIdx.x + threadIdx.x) * ILP; 
int row = blockDim.y * blockIdx.y + threadIdx.y; 
int index = row * blockDim.x * BLOCK_COUNT_X_IMPR + col; 

당신의 blockDim.x = 4 BLOCK_COUNT_X_IMPR = 100 그래서 각 행은 400 개의 인덱스를 뛰어 넘을 것입니다. 그러나, 당신의 안부가 높은 8 * 400으로

이 고려 갈 수

blockIdx = (12, 0) 
threadIdx = (3, 0) 
=> col = (12*4 + 3) * 8 = 408 
    row = 0 
    index = 408 

blockIdx = (0, 0) 
threadIdx = (1, 1) 
=> col = (0*4 + 1) * 8 = 8 
    row = 1 
    index = 1 * 400 + 8 = 408 

그래서 내가, 당신이 당신의 인덱스 작동

// gridDim.x = BLOCK_COUNT_X_IMPR 
int index = row * blockDim.x * gridDim.x * ILP + col;