使用 pthread 进行 cuda 编程

Question

#include <stdio.h>
#include <stdlib.h>
#include <pthread.h>

#define ARR_SIZE    10
#define NUM_DEVICE  1

typedef struct {
   int *arr;
   int *dev_arr;
   int *dev_result;
   int *result;
   int num;
} cuda_st;

__global__ void kernel_fc(int *dev_arr, int *dev_result)
{
    int idx = threadIdx.x;
    printf("dev_arr[%d] = %d\n", idx, dev_arr[idx]);
    atomicAdd(dev_result, dev_arr[idx]);
}

void *thread_func(void* struc)
{
    cuda_st * data = (cuda_st*)struc;
    printf("thread %d func start\n", data->num);
    printf("arr %d = ", data->num);
    for(int i=0; i<10; i++) {
        printf("%d ", data->arr[i]);
    }
    printf("\n");
    cudaSetDevice(data->num);
    cudaMemcpy(data->dev_arr, data->arr,  sizeof(int)*ARR_SIZE, cudaMemcpyHostToDevice);
    kernel_fc<<<1,ARR_SIZE>>>(data->dev_arr, data->dev_result);
    cudaMemcpy(data->result, data->dev_result, sizeof(int), cudaMemcpyDeviceToHost);
    printf("thread %d func exit\n", data->num);
    return NULL;
}

int main(void)
{
    // Make object
    cuda_st cuda[NUM_DEVICE];

    // Make thread
    pthread_t pthread[NUM_DEVICE];

    // Host array memory allocation
    int *arr[NUM_DEVICE];
    for(int i=0; i<NUM_DEVICE; i++) {
        arr[i] = (int*)malloc(sizeof(int)*ARR_SIZE);
    }

    // Fill this host array up with specified data
    for(int i=0; i<NUM_DEVICE; i++) {
        for(int j=0; j<ARR_SIZE; j++) {
            arr[i][j] = i*ARR_SIZE+j;
        }
    }

    // To confirm host array data
    for(int i=0; i<NUM_DEVICE; i++) {
        printf("arr[%d] = ", i);
        for(int j=0; j<ARR_SIZE; j++) {
            printf("%d ", arr[i][j]);
        }
        printf("\n");
    }

    // Result memory allocation
    int *result[NUM_DEVICE];
    for(int i=0; i<NUM_DEVICE; i++) {
        result[i] = (int*)malloc(sizeof(int));
        memset(result[i], 0, sizeof(int));
    }

    // Device array memory allocation
    int *dev_arr[NUM_DEVICE];
    for(int i=0; i<NUM_DEVICE; i++) {
        cudaMalloc(&dev_arr[i], sizeof(int)*ARR_SIZE);
    }

    // Device result memory allocation
    int *dev_result[NUM_DEVICE];
    for(int i=0; i<NUM_DEVICE; i++) {
        cudaMalloc(&dev_result[i], sizeof(int));
        cudaMemset(dev_result[i], 0, sizeof(int));
    }

    // Connect these pointers with object
    for(int i=0; i<NUM_DEVICE; i++) {
        cuda[i].arr = arr[i];
        cuda[i].dev_arr = dev_arr[i];
        cuda[i].result = result[i];
        cuda[i].dev_result = dev_result[i];
        cuda[i].num = i;
     }

    // Create and excute pthread
    for(int i=0; i<NUM_DEVICE; i++) {
        pthread_create(&pthread[i], NULL, thread_func, (void*)&cuda[i]);
    }

    // Join pthread
    for(int i=0; i<NUM_DEVICE; i++) {
        pthread_join(pthread[i], NULL);
    }

    for(int i=0; i<NUM_DEVICE; i++) {
        printf("result[%d] = %d\n", i, (*cuda[i].result));
    }

    return 0;
}

我制作了这样的简单测试程序，以使用多设备 cuda 代码测试 pthread。

NUM_DEVICE 设置为 1 时运行良好，设置为 2 时程序停止。

我猜是因为多个线程访问 cudaSetDevice 但我不知道如何处理。

我之前尝试过使用单主机线程和多设备（具有异步功能）来制作我的程序，但在我的情况下（不是上面的简单代码），内核函数之间有很多主机代码，因此它不能很好地异步工作.

所以我在将这种方式应用到我的真实代码之前测试在主机上使用多线程，但我遇到了这样的麻烦。

我必须在 cuda 函数和内核中使用异步函数吗？

给我一​​些建议。

Answer 1

问题是您在一台设备上分配内存。您需要在调用cudaMalloc 之前调用cudaSetDevice：

// Device array memory allocation
int *dev_arr[NUM_DEVICE];
for(int i=0; i<NUM_DEVICE; i++) {
    cudaSetDevice(i);
    cudaMalloc(&dev_arr[i], sizeof(int)*ARR_SIZE);
}

// Device result memory allocation
int *dev_result[NUM_DEVICE];
for(int i=0; i<NUM_DEVICE; i++) {
    cudaSetDevice(i);
    cudaMalloc(&dev_result[i], sizeof(int));
    cudaMemset(dev_result[i], 0, sizeof(int));
}