Sao chép một cấu trúc có chứa con trỏ đến thiết bị CUDA

Question

Tôi đang làm việc trên một dự án mà tôi cần thiết bị CUDA của mình để thực hiện tính toán trên cấu trúc chứa con trỏ.

Khi tôi cấp phát bộ nhớ cho cấu trúc và sau đó sao chép bộ nhớ vào thiết bị, nó sẽ chỉ sao chép cấu trúc chứ không phải nội dung của con trỏ. Bây giờ tôi đang làm việc xung quanh điều này bằng cách phân bổ con trỏ đầu tiên, sau đó thiết lập cấu trúc máy chủ để sử dụng con trỏ mới (mà nằm trên GPU). Mẫu mã sau đây mô tả phương pháp này sử dụng các cấu trúc từ trên cao:

#define N 10 

int main() { 

    int h_arr[N] = {1,2,3,4,5,6,7,8,9,10}; 
    StructA *h_a = (StructA*)malloc(sizeof(StructA)); 
    StructA *d_a; 
    int *d_arr; 

    // 1. Allocate device struct. 
    cudaMalloc((void**) &d_a, sizeof(StructA)); 

    // 2. Allocate device pointer. 
    cudaMalloc((void**) &(d_arr), sizeof(int)*N); 

    // 3. Copy pointer content from host to device. 
    cudaMemcpy(d_arr, h_arr, sizeof(int)*N, cudaMemcpyHostToDevice); 

    // 4. Point to device pointer in host struct. 
    h_a->arr = d_arr; 

    // 5. Copy struct from host to device. 
    cudaMemcpy(d_a, h_a, sizeof(StructA), cudaMemcpyHostToDevice); 

    // 6. Call kernel. 
    kernel<<<N,1>>>(d_a); 

    // 7. Copy struct from device to host. 
    cudaMemcpy(h_a, d_a, sizeof(StructA), cudaMemcpyDeviceToHost); 

    // 8. Copy pointer from device to host. 
    cudaMemcpy(h_arr, d_arr, sizeof(int)*N, cudaMemcpyDeviceToHost); 

    // 9. Point to host pointer in host struct. 
    h_a->arr = h_arr; 
} 

Câu hỏi của tôi là: Đây có phải là cách để làm điều đó?

Dường như có rất nhiều công việc và tôi nhắc bạn rằng đây là cấu trúc rất đơn giản. Nếu cấu trúc của tôi chứa rất nhiều con trỏ hoặc cấu trúc với con trỏ, mã để phân bổ và sao chép sẽ khá rộng và khó hiểu.

Answer 1

Edit: CUDA 6 giới thiệu Unified Memory, mà làm cho vấn đề "bản sao sâu" này dễ dàng hơn rất nhiều. Xem this post để biết thêm chi tiết.

---

Đừng quên rằng bạn có thể cấu trúc đi qua giá trị cho hạt nhân. Mã này hoạt động:

// pass struct by value (may not be efficient for complex structures) 
__global__ void kernel2(StructA in) 
{ 
    in.arr[threadIdx.x] *= 2; 
} 

Làm như vậy có nghĩa là bạn chỉ cần sao chép mảng đến thiết bị, không phải là cấu trúc:

int h_arr[N] = {1,2,3,4,5,6,7,8,9,10}; 
StructA h_a; 
int *d_arr; 

// 1. Allocate device array. 
cudaMalloc((void**) &(d_arr), sizeof(int)*N); 

// 2. Copy array contents from host to device. 
cudaMemcpy(d_arr, h_arr, sizeof(int)*N, cudaMemcpyHostToDevice); 

// 3. Point to device pointer in host struct. 
h_a.arr = d_arr; 

// 4. Call kernel with host struct as argument 
kernel2<<<N,1>>>(h_a); 

// 5. Copy pointer from device to host. 
cudaMemcpy(h_arr, d_arr, sizeof(int)*N, cudaMemcpyDeviceToHost); 

// 6. Point to host pointer in host struct 
// (or do something else with it if this is not needed) 
h_a.arr = h_arr; 

Answer 2

cấu trúc của mảng là một cơn ác mộng trong cuda. Bạn sẽ phải sao chép từng con trỏ vào một cấu trúc mới mà thiết bị có thể sử dụng. Có lẽ bạn thay vào đó có thể sử dụng một mảng các cấu trúc? Nếu không phải cách duy nhất tôi tìm thấy là tấn công nó theo cách bạn làm, điều đó hoàn toàn không đẹp.

EDIT: kể từ khi tôi không thể đưa ra ý kiến ​​về bài đầu: Bước 9 là không cần thiết, vì bạn có thể thay đổi bước 8 và 9 vào

// 8. Copy pointer from device to host. 
cudaMemcpy(h->arr, d_arr, sizeof(int)*N, cudaMemcpyDeviceToHost); 

Answer 3

Như đã chỉ ra Mark Harris, cấu trúc có thể được thông qua bởi các giá trị để CUDA hạt nhân. Tuy nhiên, một số cần được dành để thiết lập một destructor thích hợp kể từ khi destructor được gọi là lúc thoát khỏi hạt nhân.

Hãy xem xét ví dụ sau

#include <stdio.h> 

#include "Utilities.cuh" 

#define NUMBLOCKS 512 
#define NUMTHREADS 512 * 2 

/***************/ 
/* TEST STRUCT */ 
/***************/ 
struct Lock { 

    int *d_state; 

    // --- Constructor 
    Lock(void) { 
     int h_state = 0;          // --- Host side lock state initializer 
     gpuErrchk(cudaMalloc((void **)&d_state, sizeof(int))); // --- Allocate device side lock state 
     gpuErrchk(cudaMemcpy(d_state, &h_state, sizeof(int), cudaMemcpyHostToDevice)); // --- Initialize device side lock state 
    } 

    // --- Destructor (wrong version) 
    //~Lock(void) { 
    // printf("Calling destructor\n"); 
    // gpuErrchk(cudaFree(d_state)); 
    //} 

    // --- Destructor (correct version) 
// __host__ __device__ ~Lock(void) { 
//#if !defined(__CUDACC__) 
//  gpuErrchk(cudaFree(d_state)); 
//#else 
// 
//#endif 
// } 

    // --- Lock function 
    __device__ void lock(void) { while (atomicCAS(d_state, 0, 1) != 0); } 

    // --- Unlock function 
    __device__ void unlock(void) { atomicExch(d_state, 0); } 
}; 

/**********************************/ 
/* BLOCK COUNTER KERNEL WITH LOCK */ 
/**********************************/ 
__global__ void blockCounterLocked(Lock lock, int *nblocks) { 

    if (threadIdx.x == 0) { 
     lock.lock(); 
     *nblocks = *nblocks + 1; 
     lock.unlock(); 
    } 
} 

/********/ 
/* MAIN */ 
/********/ 
int main(){ 

    int h_counting, *d_counting; 
    Lock lock; 

    gpuErrchk(cudaMalloc(&d_counting, sizeof(int))); 

    // --- Locked case 
    h_counting = 0; 
    gpuErrchk(cudaMemcpy(d_counting, &h_counting, sizeof(int), cudaMemcpyHostToDevice)); 

    blockCounterLocked << <NUMBLOCKS, NUMTHREADS >> >(lock, d_counting); 
    gpuErrchk(cudaPeekAtLastError()); 
    gpuErrchk(cudaDeviceSynchronize()); 

    gpuErrchk(cudaMemcpy(&h_counting, d_counting, sizeof(int), cudaMemcpyDeviceToHost)); 
    printf("Counting in the locked case: %i\n", h_counting); 

    gpuErrchk(cudaFree(d_counting)); 
} 

với destructor không chú thích (không phải trả quá nhiều sự chú ý vào những gì mã thực sự thực hiện). Nếu bạn chạy mã đó, bạn sẽ nhận được kết quả sau đây

Calling destructor 
Counting in the locked case: 512 
Calling destructor 
GPUassert: invalid device pointer D:/Project/passStructToKernel/passClassToKernel/Utilities.cu 37 

Sau đó, có hai cuộc gọi đến trình hủy, một lần tại lối ra hạt nhân và một lần tại lối ra chính. Thông báo lỗi có liên quan đến thực tế là, nếu các vị trí bộ nhớ được trỏ đến bởi d_state được giải phóng tại lối ra hạt nhân, chúng không thể được giải phóng nữa tại lối ra chính. Theo đó, destructor phải khác nhau đối với các máy chủ và các thiết bị thực thi. Điều này được thực hiện bởi các destructor nhận xét trong đoạn mã trên.