pointers - نسخ بنية تحتوي على مؤشرات إلى جهاز CUDA




struct device (2)

كما أشار مارك هاريس ، يمكن تمرير الهياكل من خلال القيم إلى حبات CUDA. ومع ذلك ، ينبغي تكريس بعض الحذر لإنشاء destructor السليم منذ يتم استدعاء destructor في الخروج من النواة.

خذ بعين الاعتبار المثال التالي

#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));
}

مع destructor uncommented (لا تدفع الكثير من الاهتمام على ما يفعله رمز فعلا). إذا قمت بتشغيل هذا الرمز ، فسوف تتلقى الإخراج التالي

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

ثم هناك دعوتين إلى المدمر ، مرة واحدة عند الخروج من النواة ومرة ​​واحدة عند المخرج الرئيسي. ترتبط رسالة الخطأ بحقيقة أنه إذا تم تحرير مواقع الذاكرة المشار إليها بواسطة d_state عند مخرج kernel ، فلا يمكن تحريرها عند المخرج الرئيسي. وفقا لذلك ، يجب أن يكون destructor مختلفة لعمليات الإعدام المضيف والجهاز. يتم إنجاز هذا بواسطة destructor علقت في التعليمات البرمجية أعلاه.

أعمل في مشروع أحتاج فيه جهاز CUDA لإجراء عمليات حسابية على بنية تحتوي على مؤشرات.

typedef struct StructA {
    int* arr;
} StructA;

عندما أقوم بتخصيص ذاكرة للبنية ثم نسخها إلى الجهاز ، سيتم فقط نسخ البنية وليس محتوى المؤشر. الآن أنا أعمل حول هذا من خلال تخصيص المؤشر أولاً ، ثم تعيين بنية المضيف لاستخدام هذا المؤشر الجديد (الذي يتواجد على GPU). يصف نموذج التعليمات البرمجية التالي هذا الأسلوب باستخدام البنية من أعلاه:

#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;
}

سؤالي هو: هل هذه هي الطريقة للقيام بذلك؟

يبدو أن هناك الكثير من العمل ، وأنا أذكركم أن هذا هو بنية بسيطة للغاية. إذا احتوت تركيبي على الكثير من المؤشرات أو البنيات ذات المؤشرات نفسها ، فسيكون رمز التخصيص والنسخ واسعًا ومربكًا.


هيكل المصفوفات هو كابوس في كودا. سيكون عليك نسخ كل مؤشر إلى بنية جديدة يمكن للجهاز استخدامها. ربما يمكنك بدلا من ذلك استخدام مجموعة من البنى؟ إن لم تكن الطريقة الوحيدة التي وجدتها هي مهاجمتها بالطريقة التي تقوم بها ، وهي ليست بأي حال من الأحوال جميلة.

تعديل: نظرًا لأنني لا أستطيع تقديم تعليقات على المشاركة العليا: الخطوة رقم 9 زائدة عن الحاجة ، حيث يمكنك تغيير الخطوتين 8 و 9 إلى

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




host