تخفيض CUDA للعديد من المصفوفات الصغيرة غير المتكافئة الحجم

Question

أتساءل عما إذا كان بإمكان أي شخص أن يقترح أفضل طريقة لحساب المتوسط/الانحراف المعياري لعدد كبير من العناصر الصغيرة نسبيًا ولكن مصفوفات مختلفة الحجم في كودا؟

يعمل مثال التخفيض المتوازي في SDK على مصفوفة واحدة كبيرة جدًا ويبدو أن الحجم مناسب لعدد سلاسل العمليات لكل كتلة، لكن حالتي مختلفة نوعًا ما:

من الناحية النظرية، لدي عدد كبير من الكائنات التي يحتوي كل منها على مكونين، upper و lower ولكل من هذه المكونات x و أ y تنسيق.أي.

upper.x, lower.x, upper.y, lower.y

كل من هذه المصفوفات تقريبًا 800 في الطول ولكنه يختلف بين الكائنات (وليس داخل الكائن) على سبيل المثال.

Object1.lower.x = 1.1, 2.2, 3.3
Object1.lower.y = 4.4, 5.5, 6.6
Object1.upper.x = 7.7, 8.8, 9.9
Object1.upper.y = 1.1, 2.2, 3.3

Object2.lower.x = 1.0,  2.0,  3.0,  4.0, 5.0 
Object2.lower.y = 6.0,  7.0,  8.0,  9.0, 10.0
Object2.upper.x = 11.0, 12.0, 13.0, 14.0, 15.0 
Object2.upper.y = 16.0, 17.0, 18.0, 19.0, 20.0

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

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

شكرا لك مقدما

Answer 1

إذا فهمت بشكل صحيح، فأنت تريد تقليل Object1.lower.x إلى نتيجة واحدة، وObject1.lower.y إلى نتيجة أخرى، وهكذا.بالنسبة لأي كائن معين، هناك أربع مصفوفات يجب تقليلها، كلها متساوية الطول (للكائن).

هناك العديد من الطرق الممكنة لتحقيق ذلك، أحد العوامل المؤثرة هو العدد الإجمالي للكائنات في نظامك.سأفترض أن هذا الرقم كبير.

للحصول على أفضل أداء، فأنت تريد نمطًا مثاليًا للوصول إلى الذاكرة وتريد تجنب الاختلاف.نظرًا لأن عدد المصفوفات المتطابقة هو أربعة، إذا اتبعت النهج الساذج المتمثل في القيام بمصفوفة واحدة لكل مؤشر ترابط، أدناه، فلن تعاني فقط من ضعف الوصول إلى الذاكرة ولكن أيضًا سيحتاج h/w إلى التحقق من كل تكرار أي المواضيع الموجودة فيه يحتاج الالتواء إلى تنفيذ الحلقة - سيتم تعطيل تلك التي لا تفعل ذلك مما قد يكون غير فعال (خاصة إذا كانت إحدى المصفوفات أطول بكثير من المصفوفات الأخرى، على سبيل المثال).

for (int i = 0 ; i < myarraylength ; i++)
    sum += myarray[i];

بدلاً من ذلك، إذا حصلت على كل التفاف لمجموع مصفوفة واحدة، فلن يكون ذلك أكثر كفاءة فحسب، بل سيكون نمط الوصول إلى الذاكرة الخاص بك أيضًا أفضل بكثير نظرًا لأن الخيوط المجاورة ستقرأ العناصر المجاورة [1].

for (int i = tidwithinwarp ; i < warparraylength ; i += warpsize)
{
    mysum += warparray[i];
}
mysum = warpreduce(mysum);

يجب عليك أيضًا أن تأخذ محاذاة المصفوفات في الاعتبار، ويفضل أن تكون محاذاة على حد 64 بايت، على الرغم من أنه إذا كنت تقوم بالتطوير لقدرة الحوسبة 1.2 أو أعلى، فإن هذا ليس بنفس أهمية وحدات معالجة الرسومات الأقدم.

في هذا المثال، يمكنك إطلاق أربعة أعوجاج لكل كتلة، أي.128 موضوعًا، والعديد من الكتل التي تحتوي على كائنات.

[1] أنت تقول أنه يمكنك اختيار أي ترتيب للذاكرة تريده، وغالبًا ما يكون من المفيد تشذير المصفوفات بحيث تكون المصفوفة [0] [0] بجوار المصفوفة [1] [0] لأن هذا يعني أن الخيوط المجاورة يمكن أن تعمل على صفائف متجاورة والحصول على وصول مجمعة.ومع ذلك، نظرًا لأن طول المصفوفات ليس ثابتًا، فمن المحتمل أن يكون هذا الأمر معقدًا، ويتطلب حشو المصفوفات الأقصر.

Answer 2

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

هنا مثال عملي:

#include <cub/cub.cuh>
#include <cuda.h>

#include "Utilities.cuh"

#include <iostream>

#define WARPSIZE    32
#define BLOCKSIZE   256

const int N = 1024;

/*************************/
/* WARP REDUCTION KERNEL */
/*************************/
__global__ void sum(const float * __restrict__ indata, float * __restrict__ outdata) {

    unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;

    unsigned int warp_id = threadIdx.x / WARPSIZE;

    // --- Specialize WarpReduce for type float. 
    typedef cub::WarpReduce<float, WARPSIZE> WarpReduce;

    // --- Allocate WarpReduce shared memory for (N / WARPSIZE) warps
    __shared__ typename WarpReduce::TempStorage temp_storage[BLOCKSIZE / WARPSIZE];

    float result;
    if(tid < N) result = WarpReduce(temp_storage[warp_id]).Sum(indata[tid]);

    if(tid % WARPSIZE == 0) outdata[tid / WARPSIZE] = result;
}

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

    // --- Allocate host side space for 
    float *h_data       = (float *)malloc(N * sizeof(float));
    float *h_result     = (float *)malloc((N / WARPSIZE) * sizeof(float));

    float *d_data;      gpuErrchk(cudaMalloc(&d_data, N * sizeof(float)));
    float *d_result;    gpuErrchk(cudaMalloc(&d_result, (N / WARPSIZE) * sizeof(float)));

    for (int i = 0; i < N; i++) h_data[i] = (float)i;

    gpuErrchk(cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice));

    sum<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_data, d_result);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());

    gpuErrchk(cudaMemcpy(h_result, d_result, (N / WARPSIZE) * sizeof(float), cudaMemcpyDeviceToHost));

    std::cout << "output: ";
    for(int i = 0; i < (N / WARPSIZE); i++) std::cout << h_result[i] << " ";
    std::cout << std::endl;

    gpuErrchk(cudaFree(d_data));
    gpuErrchk(cudaFree(d_result));

    return 0;
}

في هذا المثال، مجموعة من الطول N يتم إنشاؤه والنتيجة هي مجموع 32 عناصر متتالية.لذا

result[0] = data[0] + ... + data[31];
result[1] = data[32] + ... + data[63];
....