उपयोगकर्ता लिखित कर्नेल के अंदर जोर

Question

मैं जोर देने के लिए एक नौसिखिया हूँ। मैं देखता हूं कि सभी जोर प्रस्तुतिकरण और उदाहरण केवल होस्ट कोड दिखाते हैं।

मैं जानना चाहता हूं कि क्या मैं अपने स्वयं के कर्नेल में डिवाइस_वेक्टर पास कर सकता हूं? कैसे? यदि हां, कर्नेल/डिवाइस कोड के अंदर इस पर ऑपरेशन की अनुमति क्या है?

Answer 1

यदि आप जोर से आवंटित/संसाधित डेटा का उपयोग करना चाहते हैं, तो आप आवंटित डेटा का कच्चा सूचक प्राप्त कर सकते हैं।

int * raw_ptr = thrust::raw_pointer_cast(dev_ptr); 

अगर आप कर्नेल मैं कभी नहीं की कोशिश की में जोर वैक्टर आवंटित करना चाहते लेकिन मुझे नहीं लगता काम करेंगे और यह भी अगर यह काम करता है मुझे नहीं लगता कि यह कोई लाभ प्रदान करेगा करते।

Answer 2

मैं इस प्रश्न का एक अद्यतन उत्तर देना चाहता हूं।

जोर 1.8 से शुरू हो रहा है, सीयूडीए थ्रस्ट प्राइमेटिव को thrust::seq निष्पादन नीति के साथ क्रमशः एक एकल CUDA थ्रेड (या अनुक्रमिक रूप से एकल सीपीयू थ्रेड के भीतर) चलाने के लिए जोड़ा जा सकता है। नीचे, एक उदाहरण की सूचना दी गई है।

आप एक धागा भीतर समानांतर निष्पादन चाहते हैं, तो आप जो कमी दिनचर्या है कि एक threadblock के भीतर से कहा जा सकता है प्रदान करता है CUB उपयोग करने पर विचार कर सकते हैं, बशर्ते कि आपके कार्ड गतिशील समानांतरवाद सक्षम बनाता है।

यहाँ के साथ थ्रस्ट

#include <stdio.h> 

#include <thrust/reduce.h> 
#include <thrust/execution_policy.h> 

/********************/ 
/* CUDA ERROR CHECK */ 
/********************/ 
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } 
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true) 
{ 
    if (code != cudaSuccess) 
    { 
     fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); 
     if (abort) exit(code); 
    } 
} 

__global__ void test(float *d_A, int N) { 

    float sum = thrust::reduce(thrust::seq, d_A, d_A + N); 

    printf("Device side result = %f\n", sum); 

} 

int main() { 

    const int N = 16; 

    float *h_A = (float*)malloc(N * sizeof(float)); 
    float sum = 0.f; 
    for (int i=0; i<N; i++) { 
     h_A[i] = i; 
     sum = sum + h_A[i]; 
    } 
    printf("Host side result = %f\n", sum); 

    float *d_A; gpuErrchk(cudaMalloc((void**)&d_A, N * sizeof(float))); 
    gpuErrchk(cudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice)); 

    test<<<1,1>>>(d_A, N); 

} 

Answer 3

यह मेरे पिछले जवाब के लिए एक अद्यतन है उदाहरण है।

थ्रस्ट 1.8.1 से शुरू, CUDA थ्रस्ट पुरातन एक भी CUDA धागा CUDA गतिशील समानांतरवाद शोषण के भीतर समानांतर में चलाने के लिए thrust::device निष्पादन नीति के साथ जोड़ा जा सकता है। नीचे, एक उदाहरण की सूचना दी गई है।

#include <stdio.h> 

#include <thrust/reduce.h> 
#include <thrust/execution_policy.h> 

#include "TimingGPU.cuh" 
#include "Utilities.cuh" 

#define BLOCKSIZE_1D 256 
#define BLOCKSIZE_2D_X 32 
#define BLOCKSIZE_2D_Y 32 

/*************************/ 
/* TEST KERNEL FUNCTIONS */ 
/*************************/ 
__global__ void test1(const float * __restrict__ d_data, float * __restrict__ d_results, const int Nrows, const int Ncols) { 

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

    if (tid < Nrows) d_results[tid] = thrust::reduce(thrust::seq, d_data + tid * Ncols, d_data + (tid + 1) * Ncols); 

} 

__global__ void test2(const float * __restrict__ d_data, float * __restrict__ d_results, const int Nrows, const int Ncols) { 

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

    if (tid < Nrows) d_results[tid] = thrust::reduce(thrust::device, d_data + tid * Ncols, d_data + (tid + 1) * Ncols); 

} 

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

    const int Nrows = 64; 
    const int Ncols = 2048; 

    gpuErrchk(cudaFree(0)); 

// size_t DevQueue; 
// gpuErrchk(cudaDeviceGetLimit(&DevQueue, cudaLimitDevRuntimePendingLaunchCount)); 
// DevQueue *= 128; 
// gpuErrchk(cudaDeviceSetLimit(cudaLimitDevRuntimePendingLaunchCount, DevQueue)); 

    float *h_data  = (float *)malloc(Nrows * Ncols * sizeof(float)); 
    float *h_results = (float *)malloc(Nrows *   sizeof(float)); 
    float *h_results1 = (float *)malloc(Nrows *   sizeof(float)); 
    float *h_results2 = (float *)malloc(Nrows *   sizeof(float)); 
    float sum = 0.f; 
    for (int i=0; i<Nrows; i++) { 
     h_results[i] = 0.f; 
     for (int j=0; j<Ncols; j++) { 
      h_data[i*Ncols+j] = i; 
      h_results[i] = h_results[i] + h_data[i*Ncols+j]; 
     } 
    } 

    TimingGPU timerGPU; 

    float *d_data;   gpuErrchk(cudaMalloc((void**)&d_data,  Nrows * Ncols * sizeof(float))); 
    float *d_results1;  gpuErrchk(cudaMalloc((void**)&d_results1, Nrows   * sizeof(float))); 
    float *d_results2;  gpuErrchk(cudaMalloc((void**)&d_results2, Nrows   * sizeof(float))); 
    gpuErrchk(cudaMemcpy(d_data, h_data, Nrows * Ncols * sizeof(float), cudaMemcpyHostToDevice)); 

    timerGPU.StartCounter(); 
    test1<<<iDivUp(Nrows, BLOCKSIZE_1D), BLOCKSIZE_1D>>>(d_data, d_results1, Nrows, Ncols); 
    gpuErrchk(cudaPeekAtLastError()); 
    gpuErrchk(cudaDeviceSynchronize()); 
    printf("Timing approach nr. 1 = %f\n", timerGPU.GetCounter()); 

    gpuErrchk(cudaMemcpy(h_results1, d_results1, Nrows * sizeof(float), cudaMemcpyDeviceToHost)); 

    for (int i=0; i<Nrows; i++) { 
     if (h_results1[i] != h_results[i]) { 
      printf("Approach nr. 1; Error at i = %i; h_results1 = %f; h_results = %f", i, h_results1[i], h_results[i]); 
      return 0; 
     } 
    } 

    timerGPU.StartCounter(); 
    test2<<<iDivUp(Nrows, BLOCKSIZE_1D), BLOCKSIZE_1D>>>(d_data, d_results1, Nrows, Ncols); 
    gpuErrchk(cudaPeekAtLastError()); 
    gpuErrchk(cudaDeviceSynchronize()); 
    printf("Timing approach nr. 2 = %f\n", timerGPU.GetCounter()); 

    gpuErrchk(cudaMemcpy(h_results1, d_results1, Nrows * sizeof(float), cudaMemcpyDeviceToHost)); 

    for (int i=0; i<Nrows; i++) { 
     if (h_results1[i] != h_results[i]) { 
      printf("Approach nr. 2; Error at i = %i; h_results1 = %f; h_results = %f", i, h_results1[i], h_results[i]); 
      return 0; 
     } 
    } 

    printf("Test passed!\n"); 

} 

ऊपर के उदाहरण Reduce matrix rows with CUDA रूप में एक ही अर्थ में एक मैट्रिक्स की पंक्तियों की कटौती करता है, लेकिन यह ऊपर पोस्ट से अलग ढंग से किया जाता है, अर्थात्, CUDA थ्रस्ट पुरातन उपयोगकर्ता लिखा कर्नेल से सीधे फोन करके। साथ ही, उपरोक्त उदाहरण दो निष्पादन नीतियों, अर्थात् thrust::seq और thrust::device के साथ किए जाने पर समान संचालन के प्रदर्शन की तुलना करने के लिए कार्य करता है।नीचे, कुछ ग्राफ प्रदर्शन में अंतर दिखाते हैं।

प्रदर्शन एक केपलर K20c पर और एक मैक्सवेल GeForce GTX 850M पर मूल्यांकन किया गया है।