المصفوفات الديناميكية CUDA العالمية (كما في C) المخصصة لذاكرة الجهاز

Question

لذا، أحاول كتابة بعض التعليمات البرمجية التي تستخدم بنية CUDA الخاصة بـ Nvidia.لقد لاحظت أن النسخ من وإلى الجهاز كان يضر حقًا بأدائي العام، لذا أحاول الآن نقل كمية كبيرة من البيانات إلى الجهاز.

ونظرًا لاستخدام هذه البيانات في العديد من الوظائف، أود أن تكون عالمية.نعم، يمكنني تمرير المؤشرات، ولكن أود حقًا أن أعرف كيفية العمل مع العناصر العالمية في هذه الحالة.

لذلك، لدي وظائف الجهاز التي تريد الوصول إلى مجموعة مخصصة للجهاز.

من الناحية المثالية، يمكنني أن أفعل شيئًا مثل:

__device__ float* global_data;

main()
{
  cudaMalloc(global_data);
  kernel1<<<blah>>>(blah); //access global data
  kernel2<<<blah>>>(blah); //access global data again
}

ومع ذلك، لم أتمكن من معرفة كيفية إنشاء مجموعة ديناميكية.لقد توصلت إلى حل بديل من خلال إعلان المصفوفة على النحو التالي:

__device__ float global_data[REALLY_LARGE_NUMBER];

وعلى الرغم من أن ذلك لا يتطلب استدعاء cudaMalloc، إلا أنني أفضل نهج التخصيص الديناميكي.

Answer 1

شيء من هذا القبيل ربما يجب أن يعمل.

#include <algorithm>

#define NDEBUG
#define CUT_CHECK_ERROR(errorMessage) do {                                 \
        cudaThreadSynchronize();                                           \
         cudaError_t err = cudaGetLastError();                             \
         if( cudaSuccess != err) {                                         \
                     fprintf(stderr, "Cuda error: %s in file '%s' in line %i : %s.\n",    \
                                             errorMessage, __FILE__, __LINE__, cudaGetErrorString( err) );\
                     exit(EXIT_FAILURE);                                                  \
                 } } while (0)


__device__ float *devPtr;

__global__
void kernel1(float *some_neat_data)
{
    devPtr = some_neat_data;
}

__global__
void kernel2(void)
{
    devPtr[threadIdx.x] *= .3f;
}


int main(int argc, char *argv[])
{
    float* otherDevPtr;
    cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr));
    cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));

    kernel1<<<1,128>>>(otherDevPtr);
    CUT_CHECK_ERROR("kernel1");

    kernel2<<<1,128>>>();

    CUT_CHECK_ERROR("kernel2");

    return 0;
}

جرب.

Answer 2

اقض بعض الوقت في التركيز على الوثائق الوفيرة التي تقدمها NVIDIA.

من دليل البرمجة:

float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));

هذا مثال بسيط لكيفية تخصيص الذاكرة.الآن، في النواة الخاصة بك، يجب عليك قبول مؤشر لتعويم مثل هذا:

__global__
void kernel1(float *some_neat_data)
{
    some_neat_data[threadIdx.x]++;
}

__global__
void kernel2(float *potentially_that_same_neat_data)
{
    potentially_that_same_neat_data[threadIdx.x] *= 0.3f;
}

والآن يمكنك استدعائهم كما يلي:

float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));

kernel1<<<1,128>>>(devPtr);
kernel2<<<1,128>>>(devPtr);

نظرًا لأن هذه البيانات تستخدم في العديد من الوظائف ، أود أن تكون عالمية.

هناك القليل من الأسباب الوجيهة لاستخدام الكرات العالمية.وهذا بالتأكيد ليس واحدًا.سأترك الأمر كتمرين لتوسيع هذا المثال ليشمل نقل "devPtr" إلى نطاق عالمي.

يحرر:

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

لذا فإن أقرب ما يمكنني اقتراحه هو:استخدم cudaMemcpyToSymbol() لتحقيق أهدافك.ولكن، في الخلفية، فكر في أن اتباع نهج مختلف قد يكون هو الشيء الصحيح.

#include <algorithm>

__constant__ float devPtr[1024];

__global__
void kernel1(float *some_neat_data)
{
    some_neat_data[threadIdx.x] = devPtr[0] * devPtr[1];
}

__global__
void kernel2(float *potentially_that_same_neat_data)
{
    potentially_that_same_neat_data[threadIdx.x] *= devPtr[2];
}


int main(int argc, char *argv[])
{
    float some_data[256];
    for (int i = 0; i < sizeof(some_data) / sizeof(some_data[0]); i++)
    {
        some_data[i] = i * 2;
    }
    cudaMemcpyToSymbol(devPtr, some_data, std::min(sizeof(some_data), sizeof(devPtr) ));
    float* otherDevPtr;
    cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr));
    cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));

    kernel1<<<1,128>>>(otherDevPtr);
    kernel2<<<1,128>>>(otherDevPtr);

    return 0;
}

لا تنسَ "--host-compilation=c++" في هذا المثال.

Answer 3

لقد تقدمت وجربت الحل المتمثل في تخصيص مؤشر مؤقت وتمريره إلى وظيفة عامة بسيطة مشابهة لـ kernel1.

والخبر السار هو أنه يعمل :)

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

على أي حال، للعلم - لقد ألقيت نظرة على العديد من الأمثلة وقمت بإجراء تمارين nvidia حيث الهدف هو الحصول على الإخراج ليقول "صحيح!".ومع ذلك، لم أنظر الجميع منهم.إذا كان أي شخص يعرف مثال sdk حيث يقومون بتخصيص ذاكرة الجهاز العالمية الديناميكية، فما زلت أرغب في معرفة ذلك.

Answer 4

حسنًا، لقد كانت مشكلة نقل devPtr إلى النطاق العالمي هي مشكلتي بالضبط.

لدي تطبيق يفعل ذلك بالضبط، حيث تحتوي النواتان على مؤشر للبيانات التي تم تمريرها.أنا صراحة لا أريد تمرير تلك المؤشرات.

لقد قرأت الوثائق عن كثب إلى حد ما، وتوجهت إلى منتديات nvidia (وبحثت في Google لمدة ساعة أو نحو ذلك)، لكنني لم أجد تطبيقًا لمصفوفة أجهزة ديناميكية عامة تعمل فعليًا (لقد جربت العديد من تلك المصفوفة والتي تجمع و ثم تفشل بطرق جديدة ومثيرة للاهتمام).

Answer 5

تحقق من العينات المضمنة مع SDK.العديد من هذه المشاريع النموذجية هي طريقة جيدة للتعلم بالقدوة.

Answer 6

ونظرًا لاستخدام هذه البيانات في العديد من الوظائف، أود أن تكون عالمية.

-

هناك القليل من الأسباب الوجيهة لاستخدام الكرات العالمية.وهذا بالتأكيد ليس واحدًا.سأترك الأمر كتمرين لتوسيع هذا المثال ليشمل نقل "DevPtr" إلى نطاق عالمي.

ماذا لو كانت النواة تعمل على بنية ثابتة كبيرة تتكون من صفائف؟إن استخدام ما يسمى بالذاكرة الثابتة ليس خيارًا، لأنها محدودة الحجم جدًا.إذن عليك أن تضعها في الذاكرة العالمية..؟