Sản phẩm CUDA Dot

Question

Tôi đang cố gắng triển khai hạt nhân dot-sản phẩm cổ điển cho các mảng chính xác kép với tính toán nguyên tử của tổng cuối cùng trên các khối khác nhau. Tôi đã sử dụng atomicAdd cho độ chính xác gấp đôi như đã nêu trong trang 116 của hướng dẫn lập trình. Có lẽ tôi đang làm điều gì đó sai. Khoản tiền một phần trên các chủ đề trong mỗi khối được tính toán một cách chính xác nhưng sau đó hoạt động nguyên tử dường như không hoạt động đúng cách vì mỗi khi tôi chạy hạt nhân với cùng một dữ liệu, tôi nhận được các kết quả khác nhau. Tôi sẽ biết ơn nếu ai đó có thể phát hiện ra sai lầm hoặc cung cấp một giải pháp thay thế! Đây là hạt nhân của tôi:

__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]); 
    } 
} 

Và đây là chức năng điện thoại của tôi 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; 
} 

Answer 1

Bạn đang sử dụng chức năng cuda_atomicAdd không chính xác. Phần này của hạt nhân của bạn:

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

là thủ phạm. Tại đây, bạn thêm nguyên tử vào dot_res. sau đó không nguyên bản đặt dot_res với kết quả trả về. Kết quả trả về từ hàm này là giá trị trước đó của vị trí đang được cập nhật nguyên tử và được cung cấp cho "thông tin" hoặc chỉ sử dụng cục bộ của người gọi. Bạn không gán nó cho những gì bạn đang cập nhật nguyên tử, mà hoàn toàn đánh bại mục đích của việc sử dụng truy cập bộ nhớ nguyên tử ngay từ đầu.Thực hiện một việc tương tự như thế này:

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

Answer 2

Bắt quyền giảm sử dụng quảng cáo hoc đang CUDA có thể được khôn lanh, vì vậy đây là một giải pháp thay thế sử dụng một thuật toán Thrust , được bao gồm trong Bộ công cụ 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); 
} 

Answer 3

Không kiểm tra mã của bạn sâu nhưng đây là một số lời khuyên.
Tôi chỉ khuyên bạn nên sử dụng Thrust nếu bạn chỉ sử dụng GPU cho các nhiệm vụ chung như vậy, vì nếu một vấn đề phức tạp phát sinh, mọi người không có ý tưởng lập trình song song hiệu quả trên gpu.

1. Bắt đầu hạt nhân giảm song song mới để tóm tắt sản phẩm chấm.
   Vì dữ liệu đã có trên thiết bị, bạn sẽ không thấy hiệu suất giảm từ hạt nhân mới.

2. Hạt nhân của bạn dường như không mở rộng trên số lượng tối đa các khối có thể có trên GPU mới nhất. Nếu nó sẽ và hạt nhân của bạn sẽ có thể tính toán các sản phẩm dấu chấm của hàng triệu giá trị hiệu suất sẽ giảm đáng kể vì hoạt động nguyên tử được tuần tự hóa.

3. Lỗi sơ cấp: Dữ liệu đầu vào của bạn và quyền truy cập bộ nhớ dùng chung có được kiểm tra không? Hoặc bạn có chắc dữ liệu đầu vào luôn là bội số của kích thước khối của bạn không? Khác bạn sẽ đọc rác. Hầu hết các kết quả sai của tôi là do lỗi này.

4. tối ưu hóa giảm song song của bạn. My Thesis hoặc Optimisations Mark Harris

chưa được kiểm tra, tôi chỉ viết nó xuống bằng notepad:

/* 
* @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! 
} 

Edit: loại bỏ các điểm không cần thiết.