CUDA求任意長度向量和debug實錄

症狀：點選運作，程式一直跑，卻沒有輸出

問題代碼：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
 
#include <stdio.h>
 
#define N 256 //向量長度，自行設定
#define BLOCK 128//線程塊數目，可根據硬體限制随意設定
#define BLOCKDIM 128 //線程塊内線程數目，可根據硬體限制随意設定
 
__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    while (i<N)
    {
        c[i] = a[i] + b[i];
        i += blockDim.x * blockIdx.x;
    }
}
 
int main()
{
    int* a = new int[N];
    int* b = new int[N];
    int* c = new int[N];
    int *dev_a = 0;
    int *dev_b = 0;
    int *dev_c = 0;
 
    for (int i = 1; i <= N; ++i) {
        a[i-1] = i;
        b[i-1] = i & 1;
    }
    cudaMalloc((void**)&dev_c, N * sizeof(int));
    cudaMalloc((void**)&dev_a, N * sizeof(int));
    cudaMalloc((void**)&dev_b, N * sizeof(int));
    cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);
 
    addKernel<<<BLOCK, BLOCKDIM>>>(dev_c, dev_a, dev_b);
    cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);
    
    for (int i = 0; i < N; ++i) {
        printf("第%d維:%d+%d=%d\n", i+1, a[i], b[i], c[i]);
    }
 
    delete[] a;
    delete[] b;
    delete[] c;
    cudaFree(dev_c);
    cudaFree(dev_a);
    cudaFree(dev_b);
 
    return 0;
}
           

第一步：在main函數中各個可疑程式段後添加printf語句，運作檢視是哪段代碼的問題，修改後的main函數

int main()
{
    int* a = new int[N];
    int* b = new int[N];
    int* c = new int[N];
    int* dev_a = 0;
    int* dev_b = 0;
    int* dev_c = 0;

    for (int i = 1; i <= N; ++i) {
        a[i - 1] = i;
        b[i - 1] = i & 1;
    }
    cudaMalloc((void**)&dev_c, N * sizeof(int));
    cudaMalloc((void**)&dev_a, N * sizeof(int));
    cudaMalloc((void**)&dev_b, N * sizeof(int));
    cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);
    printf("111111111111111\n");

    addKernel << <BLOCK, BLOCKDIM >> > (dev_c, dev_a, dev_b);
    printf("222222222222222\n");
    cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);
    printf("333333333333333\n");

    for (int i = 0; i < N; ++i) {
        printf("第%d維:%d+%d=%d\n", i + 1, a[i], b[i], c[i]);
    }

    delete[] a;
    delete[] b;
    delete[] c;
    cudaFree(dev_c);
    cudaFree(dev_a);
    cudaFree(dev_b);

    return 0;
}
           

運作修改main函數後的代碼，程式輸出

111111111111111
222222222222222
           

說明問題出在語句cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);   上網查詢cudaMemcpyDeviceToHost模式運作失敗的原因後，我嘗試了重新開機電腦，沒有用，說明是程式代碼有問題。這個模式運作失敗可能的原因還有并行處理資料的時候有部分資料超過了記憶體的邊界，會導緻處理結果無法拷貝出來。于是我仔細地檢查了add函數

__global__ void addKernel(int* c, const int* a, const int* b)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    while (i < N)
    {
        c[i] = a[i] + b[i];
        i += blockDim.x * blockIdx.x;
    }
}
           

add函數中我已經對邊界做了判斷（i<N），因而不可能是越界的問題。但是既然無法把并行處理結果拷貝出來，肯定是并行處理的過程有問題。

第二步：在add中增加printf語句，運作程式檢視各個線程塊的運作情況，修改後的add函數如下

__global__ void addKernel(int* c, const int* a, const int* b)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    while (i < N)
    {
        c[i] = a[i] + b[i];
        i += blockDim.x * blockIdx.x;
        printf("%d号線程塊", blockIdx.x);
    }
}
           

運作修改add函數後的代碼，程式瘋狂輸出“0号線程塊”，而且沒有要停止的迹象，是以可以判斷，0号線程塊出現了死循環。為什麼會這樣呢？仔細看語句i += blockDim.x * blockIdx.x;   0号block的blockIdx.x是0，是以在0号block中的每個thread的循環變量i每次增量都為0，永遠不會破壞i<N這個條件，while循環就一直進行。

這個程式的意思是重複利用固定數量的線程運算矢量的若幹次元求和。實際上，GPU一次同時運作的線程數目，是一個grid裡面橫着的一排，是以一次同時計算的矢量次元序号範圍是grid裡一橫排包括的序号範圍，是以對于每個重複利用的線程，下次計算的次元序号和這次相差一橫排的寬度。

thread、block、grid之間到底是什麼關系？詳情可以看這篇部落格的圖畫。總結來說，就是一個grid裡包含若幹block，一個block裡包含若幹thread，blockDim是一個線程格在橫排上的寬度，即橫排上分布多少個thread，gridDim是一個grid在橫排上的寬度，即橫排上分布多少個block。

經過以上分析，add代碼應修正為

__global__ void addKernel(int* c, const int* a, const int* b)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    while (i < N)
    {
        c[i] = a[i] + b[i];
        i += blockDim.x * gridDim.x;
    }
}
           

運作修正後代碼，程式正常運作。