cuda
समानांतर कमी (उदाहरण के लिए एक सरणी को कैसे योग करें)

algorithm C Language C# Language C++ Java Language GNU/Linux MATLAB Language opencv Python Language tensorflow
खोज…


टिप्पणियों

समानांतर कमी एल्गोरिथ्म आमतौर पर एक एल्गोरिथ्म को संदर्भित करता है जो तत्वों की एक सरणी को जोड़ती है, एक परिणाम उत्पन्न करता है। इस श्रेणी में आने वाली विशिष्ट समस्याएं हैं:

  • किसी सरणी में सभी तत्वों को समेटें
  • एक सरणी में अधिकतम ढूँढना

सामान्य तौर पर, समानांतर कमी किसी भी बाइनरी एसोसिएटिव ऑपरेटर के लिए लागू की जा सकती है, अर्थात (A*B)*C = A*(B*C) । ऐसे ऑपरेटर * के साथ, समानांतर कमी एल्गोरिदम दोहराए गए समूहों को जोड़ियों में सरणी तर्क देता है। प्रत्येक जोड़ी को एक चरण में समग्र सरणी आकार को आधा करते हुए, दूसरों के साथ समानांतर में गणना की जाती है। प्रक्रिया तब तक दोहराई जाती है जब तक केवल एक ही तत्व मौजूद न हो।

यदि ऑपरेटर कम्यूटेटिव है (यानी A*B = B*A ) सहयोगी होने के अलावा, एल्गोरिथ्म एक अलग पैटर्न में जोड़ सकता है। सैद्धांतिक दृष्टिकोण से इसे कोई फर्क नहीं पड़ता है, लेकिन व्यवहार में यह बेहतर मेमोरी एक्सेस पैटर्न देता है:

सभी सहयोगी ऑपरेटर्स कम्यूटेटिव नहीं हैं - उदाहरण के लिए मैट्रिक्स गुणन लें।

कम्यूटेटर ऑपरेटर के लिए एकल-ब्लॉक समानांतर कमी

CUDA में समानांतर कमी के लिए सबसे सरल तरीका है, कार्य करने के लिए एक ही ब्लॉक को असाइन करना: 

static const int arraySize = 10000;
static const int blockSize = 1024;

__global__ void sumCommSingleBlock(const int *a, int *out) {
    int idx = threadIdx.x;
    int sum = 0;
    for (int i = idx; i < arraySize; i += blockSize)
        sum += a[i];
    __shared__ int r[blockSize];
    r[idx] = sum;
    __syncthreads();
    for (int size = blockSize/2; size>0; size/=2) { //uniform
        if (idx<size)
            r[idx] += r[idx+size];
        __syncthreads();
    }
    if (idx == 0)
        *out = r[0];
}

...

sumCommSingleBlock<<<1, blockSize>>>(dev_a, dev_out);

यह सबसे अधिक संभव है जब डेटा का आकार बहुत बड़ा नहीं होता है (कुछ कंप्रेशेंट तत्वों के आसपास)। यह आमतौर पर तब होता है जब कटौती कुछ बड़े CUDA कार्यक्रम का एक हिस्सा है। इनपुट से मेल खाता है blockSize से बहुत शुरुआत, पहले for पाश पूरी तरह से हटाया जा सकता है।

ध्यान दें कि पहले चरण में, जब थ्रेड से अधिक तत्व होते हैं, तो हम चीजों को पूरी तरह से स्वतंत्र रूप से जोड़ते हैं। केवल तभी जब समस्या को कम किया blockSize , वास्तविक समानांतर कमी ट्रिगर होती है। समान कोड किसी भी अन्य कम्यूटेटिव, एसोसिएटिव ऑपरेटर पर लागू किया जा सकता है, जैसे गुणा, न्यूनतम, अधिकतम, आदि।

ध्यान दें कि एल्गोरिथ्म तेजी से बनाया जा सकता है, उदाहरण के लिए एक ताना-स्तर समानांतर कमी का उपयोग करके।

गैर-कम्यूटेटिव ऑपरेटर के लिए एकल-ब्लॉक समानांतर कमी

एक गैर-कम्यूटेटिव ऑपरेटर के लिए समानांतर कमी करना कम्यूटेटिव संस्करण की तुलना में थोड़ा अधिक शामिल है। उदाहरण में हम अभी भी सादगी के लिए पूर्णांक पर एक अतिरिक्त का उपयोग करते हैं। यह प्रतिस्थापित किया जा सकता है, उदाहरण के लिए, मैट्रिक्स गुणा के साथ जो वास्तव में गैर-कम्यूटेटिव है। ध्यान दें, ऐसा करते समय, 0 को गुणन के एक तटस्थ तत्व द्वारा प्रतिस्थापित किया जाना चाहिए - अर्थात एक पहचान मैट्रिक्स। 

static const int arraySize = 1000000;
static const int blockSize = 1024;

__global__ void sumNoncommSingleBlock(const int *gArr, int *out) {
    int thIdx = threadIdx.x;
    __shared__ int shArr[blockSize*2];
    __shared__ int offset;
    shArr[thIdx] = thIdx<arraySize ? gArr[thIdx] : 0;
    if (thIdx == 0)
        offset = blockSize;
    __syncthreads();
    while (offset < arraySize) { //uniform
        shArr[thIdx + blockSize] = thIdx+offset<arraySize ? gArr[thIdx+offset] : 0;
        __syncthreads();
        if (thIdx == 0)
            offset += blockSize;
        int sum = shArr[2*thIdx] + shArr[2*thIdx+1];
        __syncthreads();
        shArr[thIdx] = sum;
    }
    __syncthreads();
    for (int stride = 1; stride<blockSize; stride*=2) { //uniform
        int arrIdx = thIdx*stride*2;
        if (arrIdx+stride<blockSize)
            shArr[arrIdx] += shArr[arrIdx+stride];
        __syncthreads();
    }   
    if (thIdx == 0)
        *out = shArr[0];
}

...

sumNoncommSingleBlock<<<1, blockSize>>>(dev_a, dev_out);

पहले जबकि लूप निष्पादित होता है जब तक कि थ्रेड की तुलना में अधिक इनपुट तत्व होते हैं। प्रत्येक पुनरावृत्ति में, एक एकल कमी की जाती है और परिणाम को shArr सरणी के पहले आधे भाग में संपीड़ित किया जाता है। दूसरी छमाही फिर नए डेटा से भर जाती है।

एक बार सभी डेटा gArr से लोड होने के gArr , दूसरा लूप निष्पादित होता है। अब, हम अब परिणाम को संक्षिप्त नहीं करते हैं (जिसकी अतिरिक्त __syncthreads() ) लागत होती है। प्रत्येक चरण में थ्रेड n 2*n -th सक्रिय तत्व तक पहुँचता है और इसे 2*n+1 -थ तत्व के साथ जोड़ता है:

इस सरल उदाहरण को आगे बढ़ाने के कई तरीके हैं, उदाहरण के लिए ताना-स्तर में कमी और साझा मेमोरी बैंक संघर्षों को हटाकर।

कम्यूटेटर ऑपरेटर के लिए बहु-ब्लॉक समानांतर कमी

एकल-ब्लॉक दृष्टिकोण की तुलना में CUDA में समानांतर कमी के लिए मल्टी-ब्लॉक दृष्टिकोण एक अतिरिक्त चुनौती बन गया है, क्योंकि ब्लॉक संचार में सीमित हैं। विचार यह है कि प्रत्येक ब्लॉक को इनपुट सरणी के एक भाग की गणना करने दें, और फिर सभी आंशिक परिणामों को मर्ज करने के लिए एक अंतिम ब्लॉक है। ऐसा करने के लिए, एक दो गुठली को लॉन्च कर सकता है, जिसका अर्थ है कि एक ग्रिड-वाइड सिंक्रनाइज़ेशन बिंदु बनाना। 

static const int wholeArraySize = 100000000;
static const int blockSize = 1024;
static const int gridSize = 24; //this number is hardware-dependent; usually #SM*2 is a good number.

__global__ void sumCommMultiBlock(const int *gArr, int arraySize, int *gOut) {
    int thIdx = threadIdx.x;
    int gthIdx = thIdx + blockIdx.x*blockSize;
    const int gridSize = blockSize*gridDim.x;
    int sum = 0;
    for (int i = gthIdx; i < arraySize; i += gridSize)
        sum += gArr[i];
    __shared__ int shArr[blockSize];
    shArr[thIdx] = sum;
    __syncthreads();
    for (int size = blockSize/2; size>0; size/=2) { //uniform
        if (thIdx<size)
            shArr[thIdx] += shArr[thIdx+size];
        __syncthreads();
    }
    if (thIdx == 0)
        gOut[blockIdx.x] = shArr[0];
}

__host__ int sumArray(int* arr) {
    int* dev_arr;
    cudaMalloc((void**)&dev_arr, wholeArraySize * sizeof(int));
    cudaMemcpy(dev_arr, arr, wholeArraySize * sizeof(int), cudaMemcpyHostToDevice);

    int out;
    int* dev_out;
    cudaMalloc((void**)&dev_out, sizeof(int)*gridSize);
    
    sumCommMultiBlock<<<gridSize, blockSize>>>(dev_arr, wholeArraySize, dev_out);
    //dev_out now holds the partial result
    sumCommMultiBlock<<<1, blockSize>>>(dev_out, gridSize, dev_out);
    //dev_out[0] now holds the final result
    cudaDeviceSynchronize();
    
    cudaMemcpy(&out, dev_out, sizeof(int), cudaMemcpyDeviceToHost);
    cudaFree(dev_arr);
    cudaFree(dev_out);
    return out;
}

आदर्श रूप से पूरी व्यस्तता में GPU पर सभी मल्टीप्रोसेसर को संतृप्त करने के लिए पर्याप्त ब्लॉक लॉन्च करना चाहता है। इस संख्या से अधिक - विशेष रूप से, उतने ही धागे उतारे जा सकते हैं जितने कि सरणी में तत्व हों - प्रति-उत्पादक। ऐसा करने से यह कच्ची गणना शक्ति में वृद्धि नहीं करता है, लेकिन बहुत कुशल पहले लूप का उपयोग करने से रोकता है।

अंतिम-ब्लॉक गार्ड की सहायता से एकल कर्नेल का उपयोग करके समान परिणाम प्राप्त करना भी संभव है: 

static const int wholeArraySize = 100000000;
static const int blockSize = 1024;
static const int gridSize = 24;

__device__ bool lastBlock(int* counter) {
    __threadfence(); //ensure that partial result is visible by all blocks
    int last = 0;
    if (threadIdx.x == 0)
        last = atomicAdd(counter, 1);
    return __syncthreads_or(last == gridDim.x-1);
}    

__global__ void sumCommMultiBlock(const int *gArr, int arraySize, int *gOut, int* lastBlockCounter) {
    int thIdx = threadIdx.x;
    int gthIdx = thIdx + blockIdx.x*blockSize;
    const int gridSize = blockSize*gridDim.x;
    int sum = 0;
    for (int i = gthIdx; i < arraySize; i += gridSize)
        sum += gArr[i];
    __shared__ int shArr[blockSize];
    shArr[thIdx] = sum;
    __syncthreads();
    for (int size = blockSize/2; size>0; size/=2) { //uniform
        if (thIdx<size)
            shArr[thIdx] += shArr[thIdx+size];
        __syncthreads();
    }
    if (thIdx == 0)
        gOut[blockIdx.x] = shArr[0];
    if (lastBlock(lastBlockCounter)) {
        shArr[thIdx] = thIdx<gridSize ? gOut[thIdx] : 0;
        __syncthreads();
        for (int size = blockSize/2; size>0; size/=2) { //uniform
            if (thIdx<size)
                shArr[thIdx] += shArr[thIdx+size];
            __syncthreads();
        }
        if (thIdx == 0)
            gOut[0] = shArr[0];            
    }
}

__host__ int sumArray(int* arr) {
    int* dev_arr;
    cudaMalloc((void**)&dev_arr, wholeArraySize * sizeof(int));
    cudaMemcpy(dev_arr, arr, wholeArraySize * sizeof(int), cudaMemcpyHostToDevice);

    int out;
    int* dev_out;
    cudaMalloc((void**)&dev_out, sizeof(int)*gridSize);
    
    int* dev_lastBlockCounter;
    cudaMalloc((void**)&dev_lastBlockCounter, sizeof(int));
    cudaMemset(dev_lastBlockCounter, 0, sizeof(int));
    
    sumCommMultiBlock<<<gridSize, blockSize>>>(dev_arr, wholeArraySize, dev_out, dev_lastBlockCounter);
    cudaDeviceSynchronize();
    
    cudaMemcpy(&out, dev_out, sizeof(int), cudaMemcpyDeviceToHost);
    cudaFree(dev_arr);
    cudaFree(dev_out);
    return out;
}

ध्यान दें कि कर्नेल तेजी से बनाया जा सकता है, उदाहरण के लिए एक ताना-स्तर समानांतर कमी का उपयोग करके।

नॉन-ऑपरेटिव ऑपरेटर के लिए मल्टी-ब्लॉक समानांतर कमी

बहु-ब्लॉक दृष्टिकोण समानांतर कमी के लिए एकल-ब्लॉक दृष्टिकोण के समान है। वैश्विक इनपुट सरणी को खंडों में विभाजित किया जाना चाहिए, प्रत्येक को एक ब्लॉक द्वारा घटाया जाएगा। जब प्रत्येक ब्लॉक से आंशिक परिणाम प्राप्त होता है, तो अंतिम परिणाम प्राप्त करने के लिए एक अंतिम ब्लॉक इन्हें कम कर देता है।

  • sumNoncommSingleBlock को एकल-ब्लॉक कमी उदाहरण में अधिक विस्तार से समझाया गया है।
  • lastBlock केवल उस तक पहुंचने वाले अंतिम ब्लॉक को स्वीकार करता है। यदि आप इससे बचना चाहते हैं, तो आप कर्नेल को दो अलग-अलग कॉल में विभाजित कर सकते हैं। 
static const int wholeArraySize = 100000000;
static const int blockSize = 1024;
static const int gridSize = 24; //this number is hardware-dependent; usually #SM*2 is a good number.

__device__ bool lastBlock(int* counter) {
    __threadfence(); //ensure that partial result is visible by all blocks
    int last = 0;
    if (threadIdx.x == 0)
        last = atomicAdd(counter, 1);
    return __syncthreads_or(last == gridDim.x-1);
}

__device__ void sumNoncommSingleBlock(const int* gArr, int arraySize, int* out) {
    int thIdx = threadIdx.x;
    __shared__ int shArr[blockSize*2];
    __shared__ int offset;
    shArr[thIdx] = thIdx<arraySize ? gArr[thIdx] : 0;
    if (thIdx == 0)
        offset = blockSize;
    __syncthreads();
    while (offset < arraySize) { //uniform
        shArr[thIdx + blockSize] = thIdx+offset<arraySize ? gArr[thIdx+offset] : 0;
        __syncthreads();
        if (thIdx == 0)
            offset += blockSize;
        int sum = shArr[2*thIdx] + shArr[2*thIdx+1];
        __syncthreads();
        shArr[thIdx] = sum;
    }
    __syncthreads();
    for (int stride = 1; stride<blockSize; stride*=2) { //uniform
        int arrIdx = thIdx*stride*2;
        if (arrIdx+stride<blockSize)
            shArr[arrIdx] += shArr[arrIdx+stride];
        __syncthreads();
    }   
    if (thIdx == 0)
        *out = shArr[0];
}

__global__ void sumNoncommMultiBlock(const int* gArr, int* out, int* lastBlockCounter) {
    int arraySizePerBlock = wholeArraySize/gridSize;
    const int* gArrForBlock = gArr+blockIdx.x*arraySizePerBlock;
    int arraySize = arraySizePerBlock;
    if (blockIdx.x == gridSize-1)
        arraySize = wholeArraySize - blockIdx.x*arraySizePerBlock;
    sumNoncommSingleBlock(gArrForBlock, arraySize, &out[blockIdx.x]);
    if (lastBlock(lastBlockCounter))
        sumNoncommSingleBlock(out, gridSize, out);    
}

आदर्श रूप से पूरी व्यस्तता में GPU पर सभी मल्टीप्रोसेसर को संतृप्त करने के लिए पर्याप्त ब्लॉक लॉन्च करना चाहता है। इस संख्या से अधिक - विशेष रूप से, उतने ही धागे उतारे जा सकते हैं जितने कि सरणी में तत्व हों - प्रति-उत्पादक। ऐसा करने से यह कच्ची गणना शक्ति में वृद्धि नहीं करता है, लेकिन बहुत कुशल पहले लूप का उपयोग करने से रोकता है।

कम्यूटेटिव ऑपरेटर के लिए एकल-ताना समानांतर कमी

कभी-कभी कमी को एक बड़े CUDA कर्नेल के एक भाग के रूप में बहुत छोटे पैमाने पर प्रदर्शन करना पड़ता है। उदाहरण के लिए मान लें, कि इनपुट डेटा में 32 तत्व हैं - एक ताना में धागे की संख्या। ऐसे परिदृश्य में कमी को पूरा करने के लिए एक ही ताना दिया जा सकता है। यह देखते हुए कि ताना एक परिपूर्ण सिंक में निष्पादित होता है, कई __syncthreads() निर्देशों को हटाया जा सकता है - जब एक ब्लॉक-स्तरीय कमी की तुलना में। 

static const int warpSize = 32;

__device__ int sumCommSingleWarp(volatile int* shArr) {
    int idx = threadIdx.x % warpSize; //the lane index in the warp
    if (idx<16) shArr[idx] += shArr[idx+16];
    if (idx<8) shArr[idx] += shArr[idx+8];
    if (idx<4) shArr[idx] += shArr[idx+4];
    if (idx<2) shArr[idx] += shArr[idx+2];
    if (idx==0) shArr[idx] += shArr[idx+1];
    return shArr[0];
}

shArr साझा मेमोरी में एक सरणी है। ताना में सभी धागे के लिए मूल्य समान होना चाहिए। यदि sumCommSingleWarp को कई sumCommSingleWarp द्वारा बुलाया जाता है, तो shArr (प्रत्येक shArr समान) के बीच अलग होना चाहिए।

तर्क shArr को यह सुनिश्चित करने के लिए volatile के रूप में चिह्नित किया गया है कि सरणी पर संचालन वास्तव में जहां संकेत किया गया है। अन्यथा, shArr[idx] को shArr[idx] जाने वाले असाइनमेंट को एक रजिस्टर में असाइनमेंट के रूप में अनुकूलित किया जा सकता है, केवल अंतिम shArr लिए shArr को एक वास्तविक स्टोर होना shArr । जब ऐसा होता है, तो तत्काल कार्य अन्य थ्रेड्स के लिए दिखाई नहीं देते हैं, गलत परिणाम देते हैं। ध्यान दें, कि आप एक सामान्य गैर-वाष्पशील सरणी को अस्थिर एक के तर्क के रूप में पास कर सकते हैं, उसी तरह जब आप एक नॉन-कास्ट को एक कास्ट पैरामीटर के रूप में पास करते हैं।

यदि कोई कमी के बाद shArr[1..31] की सामग्री की परवाह नहीं करता है, तो कोई आगे भी कोड को सरल बना सकता है: 

static const int warpSize = 32;

__device__ int sumCommSingleWarp(volatile int* shArr) {
    int idx = threadIdx.x % warpSize; //the lane index in the warp
    if (idx<16) {
      shArr[idx] += shArr[idx+16];
      shArr[idx] += shArr[idx+8];
      shArr[idx] += shArr[idx+4];
      shArr[idx] += shArr[idx+2];
      shArr[idx] += shArr[idx+1];
    }
    return shArr[0];
}

इस सेटअप में हम कई हटा दिया if स्थिति। अतिरिक्त धागे कुछ अनावश्यक जोड़-घटाव करते हैं, लेकिन हम अब उन सामग्रियों की परवाह नहीं करते हैं जो वे पैदा करते हैं। चूंकि SIMD मोड में वॉर निष्पादित होते हैं, हम वास्तव में उन थ्रेड्स को कुछ भी नहीं होने से समय पर नहीं बचाते हैं। दूसरी ओर, स्थितियों का मूल्यांकन करने में अपेक्षाकृत बड़ी मात्रा में समय लगता है, क्योंकि इनमें से if विवरण इतने छोटे हैं। if प्रारंभिक कथन को हटाया जा सकता है तो shArr[32..47] 0 के साथ गद्देदार है।

ताना-स्तर में कमी का उपयोग ब्लॉक-स्तर में कमी को बढ़ावा देने के लिए भी किया जा सकता है: 

__global__ void sumCommSingleBlockWithWarps(const int *a, int *out) {
    int idx = threadIdx.x;
    int sum = 0;
    for (int i = idx; i < arraySize; i += blockSize)
        sum += a[i];
    __shared__ int r[blockSize];
    r[idx] = sum;
    sumCommSingleWarp(&r[idx & ~(warpSize-1)]);
    __syncthreads();
    if (idx<warpSize) { //first warp only
        r[idx] = idx*warpSize<blockSize ? r[idx*warpSize] : 0;
        sumCommSingleWarp(r);
        if (idx == 0)
            *out = r[0];
    }
}

तर्क &r[idx & ~(warpSize-1)] मूल रूप से r + warpIdx*32 । यह प्रभावी रूप से r सरणी को 32 तत्वों के विखंडू में विभाजित करता है, और प्रत्येक चंक को अलग-अलग ताना दिया जाता है।

गैर-संचालक ऑपरेटर के लिए एकल-ताना समानांतर कमी

कभी-कभी कमी को एक बड़े CUDA कर्नेल के एक भाग के रूप में बहुत छोटे पैमाने पर प्रदर्शन करना पड़ता है। उदाहरण के लिए मान लें, कि इनपुट डेटा में 32 तत्व हैं - एक ताना में धागे की संख्या। ऐसे परिदृश्य में कमी को पूरा करने के लिए एक ही ताना दिया जा सकता है। यह देखते हुए कि ताना एक परिपूर्ण सिंक में निष्पादित होता है, कई __syncthreads() निर्देशों को हटाया जा सकता है - जब एक ब्लॉक-स्तरीय कमी की तुलना में। 

static const int warpSize = 32;

__device__ int sumNoncommSingleWarp(volatile int* shArr) {
    int idx = threadIdx.x % warpSize; //the lane index in the warp
    if (idx%2 == 0) shArr[idx] += shArr[idx+1];
    if (idx%4 == 0) shArr[idx] += shArr[idx+2];
    if (idx%8 == 0) shArr[idx] += shArr[idx+4];
    if (idx%16 == 0) shArr[idx] += shArr[idx+8];
    if (idx == 0) shArr[idx] += shArr[idx+16];
    return shArr[0];
}

shArr साझा मेमोरी में एक सरणी है। ताना में सभी धागे के लिए मूल्य समान होना चाहिए। यदि sumCommSingleWarp को कई sumCommSingleWarp द्वारा बुलाया जाता है, तो shArr (प्रत्येक shArr समान) के बीच अलग होना चाहिए।

तर्क shArr को यह सुनिश्चित करने के लिए volatile के रूप में चिह्नित किया गया है कि सरणी पर संचालन वास्तव में जहां संकेत किया गया है। अन्यथा, shArr[idx] को shArr[idx] जाने वाले असाइनमेंट को एक रजिस्टर में असाइनमेंट के रूप में अनुकूलित किया जा सकता है, केवल अंतिम shArr लिए shArr को एक वास्तविक स्टोर होना shArr । जब ऐसा होता है, तो तत्काल कार्य अन्य थ्रेड्स के लिए दिखाई नहीं देते हैं, गलत परिणाम देते हैं। ध्यान दें, कि आप एक सामान्य गैर-वाष्पशील सरणी को अस्थिर एक के तर्क के रूप में पास कर सकते हैं, उसी तरह जब आप एक नॉन-कास्ट को एक कास्ट पैरामीटर के रूप में पास करते हैं।

एक के अंतिम सामग्री के बारे में परवाह नहीं है, तो shArr[1..31] और कर सकते हैं पैड shArr[32..47] शून्य से, एक उपरोक्त कोड को आसान बनाने में कर सकते हैं: 

static const int warpSize = 32;

__device__ int sumNoncommSingleWarpPadded(volatile int* shArr) {
    //shArr[32..47] == 0
    int idx = threadIdx.x % warpSize; //the lane index in the warp
    shArr[idx] += shArr[idx+1];
    shArr[idx] += shArr[idx+2];
    shArr[idx] += shArr[idx+4];
    shArr[idx] += shArr[idx+8];
    shArr[idx] += shArr[idx+16];
    return shArr[0];
}

इस सेटअप में हम सभी को हटा दिया if स्थिति है, जो निर्देश के आधे के बारे में बनाते हैं। अतिरिक्त थ्रेड्स कुछ अनावश्यक जोड़- shArr करते हैं, परिणाम को shArr कोशिकाओं में shArr जो अंततः अंतिम परिणाम पर कोई प्रभाव नहीं डालते हैं। चूंकि SIMD मोड में वॉर निष्पादित होते हैं, हम वास्तव में उन थ्रेड्स को कुछ भी नहीं होने से समय पर नहीं बचाते हैं।

केवल रजिस्टरों का उपयोग करके एकल-ताना समानांतर कमी

आमतौर पर, कमी वैश्विक या साझा सरणी पर की जाती है। हालाँकि, जब कमी एक बहुत छोटे पैमाने पर की जाती है, तो एक बड़े CUDA कर्नेल के एक भाग के रूप में, यह एक एकल ताना के साथ किया जा सकता है। जब ऐसा होता है, तो केप्लर या उच्चतर आर्किटेक्चर (CC> = 3.0) पर, साझा मेमोरी का उपयोग करने से बचने के लिए ताना-फेरबदल कार्यों का उपयोग करना संभव है।

उदाहरण के लिए मान लें, कि एक ताना में प्रत्येक धागा एक एकल इनपुट डेटा मान रखता है। सभी थ्रेड्स में एक साथ 32 तत्व होते हैं, जिन्हें हमें समेटने (या अन्य सहयोगी ऑपरेशन करने) की आवश्यकता होती है 

__device__ int sumSingleWarpReg(int value) {
    value += __shfl_down(value, 1);
    value += __shfl_down(value, 2);
    value += __shfl_down(value, 4);
    value += __shfl_down(value, 8);
    value += __shfl_down(value, 16);
    return __shfl(value,0);
}

यह संस्करण कम्यूटेटिव और गैर-कम्यूटेटिव ऑपरेटर्स दोनों के लिए काम करता है।



Modified text is an extract of the original Stack Overflow Documentation
के तहत लाइसेंस प्राप्त है CC BY-SA 3.0
से संबद्ध नहीं है Stack Overflow