क्या memcpy() के बराबर है जो एक CUDA कर्नेल के अंदर काम करता है?

Question

मैं अलग-अलग तोड़ने की कोशिश कर रहा हूं और CUDA कर्नेल का उपयोग करके एक सरणी के रूप में एक सरणी की संरचना को दोबारा बदल रहा हूं। memcpy() कर्नेल के अंदर काम नहीं करता है, और न ही cudaMemcpy() * करता है; मुझे हानि हो रही है।

क्या कोई मुझे CUDA कर्नेल के भीतर से स्मृति की प्रतिलिपि बनाने के लिए पसंदीदा विधि बता सकता है?

यह ध्यान देने योग्य है, cudaMemcpy(void *to, void *from, size, cudaMemcpyDeviceToDevice) जो मैं करने की कोशिश कर रहा हूं उसके लिए काम नहीं करेगा, क्योंकि इसे केवल कर्नेल के बाहर से ही बुलाया जा सकता है और यह असीमित रूप से निष्पादित नहीं होता है।

Answer 1

हां, memcpy के बराबर है जो कुडा कर्नेल के अंदर काम करता है। इसे memcpy कहा जाता है।

__global__ void kernel(int **in, int **out, int len, int N) 
{ 
    int idx = threadIdx.x + blockIdx.x*blockDim.x; 

    for(; idx<N; idx+=gridDim.x*blockDim.x) 
     memcpy(out[idx], in[idx], sizeof(int)*len); 

} 

जो इस तरह बिना किसी त्रुटि के संकलित:

$ nvcc -Xptxas="-v" -arch=sm_20 -c memcpy.cu 
ptxas info : Compiling entry function '_Z6kernelPPiS0_ii' for 'sm_20' 
ptxas info : Function properties for _Z6kernelPPiS0_ii 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 11 registers, 48 bytes cmem[0] 

और उत्सर्जन करता है PTX: एक उदाहरण के रूप

.version 3.0 
.target sm_20 
.address_size 32 

    .file 1 "/tmp/tmpxft_00000407_00000000-9_memcpy.cpp3.i" 
    .file 2 "memcpy.cu" 
    .file 3 "/usr/local/cuda/nvvm/ci_include.h" 

.entry _Z6kernelPPiS0_ii(
    .param .u32 _Z6kernelPPiS0_ii_param_0, 
    .param .u32 _Z6kernelPPiS0_ii_param_1, 
    .param .u32 _Z6kernelPPiS0_ii_param_2, 
    .param .u32 _Z6kernelPPiS0_ii_param_3 
) 
{ 
    .reg .pred %p<4>; 
    .reg .s32 %r<32>; 
    .reg .s16 %rc<2>; 


    ld.param.u32 %r15, [_Z6kernelPPiS0_ii_param_0]; 
    ld.param.u32 %r16, [_Z6kernelPPiS0_ii_param_1]; 
    ld.param.u32 %r2, [_Z6kernelPPiS0_ii_param_3]; 
    cvta.to.global.u32 %r3, %r15; 
    cvta.to.global.u32 %r4, %r16; 
    .loc 2 4 1 
    mov.u32  %r5, %ntid.x; 
    mov.u32  %r17, %ctaid.x; 
    mov.u32  %r18, %tid.x; 
    mad.lo.s32 %r30, %r5, %r17, %r18; 
    .loc 2 6 1 
    setp.ge.s32  %p1, %r30, %r2; 
    @%p1 bra BB0_5; 

    ld.param.u32 %r26, [_Z6kernelPPiS0_ii_param_2]; 
    shl.b32  %r7, %r26, 2; 
    .loc 2 6 54 
    mov.u32  %r19, %nctaid.x; 
    .loc 2 4 1 
    mov.u32  %r29, %ntid.x; 
    .loc 2 6 54 
    mul.lo.s32 %r8, %r29, %r19; 

BB0_2: 
    .loc 2 7 1 
    shl.b32  %r21, %r30, 2; 
    add.s32  %r22, %r4, %r21; 
    ld.global.u32 %r11, [%r22]; 
    add.s32  %r23, %r3, %r21; 
    ld.global.u32 %r10, [%r23]; 
    mov.u32  %r31, 0; 

BB0_3: 
    add.s32  %r24, %r10, %r31; 
    ld.u8 %rc1, [%r24]; 
    add.s32  %r25, %r11, %r31; 
    st.u8 [%r25], %rc1; 
    add.s32  %r31, %r31, 1; 
    setp.lt.u32  %p2, %r31, %r7; 
    @%p2 bra BB0_3; 

    .loc 2 6 54 
    add.s32  %r30, %r8, %r30; 
    ld.param.u32 %r27, [_Z6kernelPPiS0_ii_param_3]; 
    .loc 2 6 1 
    setp.lt.s32  %p3, %r30, %r27; 
    @%p3 bra BB0_2; 

BB0_5: 
    .loc 2 9 2 
    ret; 
} 

BB0_3 पर कोड ब्लॉक एक बाइट आकार memcpy पाश द्वारा पूर्ण रूप से अपने उत्सर्जित है कंपाइलर यह प्रदर्शन बिंदु के दृष्टिकोण से इसका एक अच्छा विचार नहीं हो सकता है, लेकिन यह पूरी तरह से समर्थित है (और सभी आर्किटेक्चर पर लंबे समय से रहा है)।

---

कि जोड़ने के लिए चार साल बाद संपादित के बाद से डिवाइस पक्ष क्रम एपीआई CUDA 6 रिलीज चक्र का हिस्सा के रूप में जारी किया गया था, यह भी संभव है सीधे डिवाइस कोड में की तरह

cudaMemcpy(void *to, void *from, size, cudaMemcpyDeviceToDevice) 

कुछ कॉल करने के लिए सभी आर्किटेक्चर जो इसका समर्थन करते हैं (गणना क्षमता 3.5 और नए हार्डवेयर) के लिए।

Answer 2

cudaMemcpy() वास्तव में असीमित रूप से चलाता है लेकिन आप सही हैं, इसे कर्नेल के भीतर से निष्पादित नहीं किया जा सकता है।

क्या कुछ गणना के आधार पर निर्धारित सरणी का नया आकार निर्धारित है? फिर, आप आमतौर पर समान संख्या में थ्रेड चलाते हैं क्योंकि आपकी सरणी में प्रविष्टियां होती हैं। प्रत्येक धागा सरणी में एक प्रविष्टि के स्रोत और गंतव्य को निर्धारित करने के लिए गणना चलाएगा और उसके बाद इसे एक असाइनमेंट के साथ कॉपी करेगा। (dst[i] = src[j])। यदि सरणी का नया आकार गणनाओं पर आधारित नहीं है, तो होस्ट से cudaMemCpyDeviceToDevice के साथ cudaMemcpy() की श्रृंखला चलाने के लिए यह अधिक कुशल हो सकता है।

Answer 3

मेरे परीक्षण में सबसे अच्छा जवाब है अपनी खुद की लूपिंग प्रति दिनचर्या लिखना। मेरे मामले में:

एक कर्नेल में

__device__ 
void devCpyCplx(const thrust::complex<float> *in, thrust::complex<float> *out, int len){ 
    // Casting for improved loads and stores 
    for (int i=0; i<len/2; ++i) { 
    ((float4*) out)[i] = ((float4*) out)[i]; 
    } 
    if (len%2) { 
    ((float2*) out)[len-1] = ((float2*) in)[len-1]; 
    } 
} 

memcpy काम करता है, लेकिन यह बहुत धीमी हो सकती है। मेजबान से cudaMemcpyAsync एक वैध विकल्प है।

मुझे 1,600 प्रतिलिपि कॉल के साथ विभिन्न बफर में ~ 33,000 लंबाई से 16,500 लंबाई के 800 संगत वैक्टरों को विभाजित करने की आवश्यकता थी।

• कर्नेल में memcpy: कर्नेल में 34 एमएस
• पाश प्रतिलिपि: 140 एमएस
• cudaMemcpy मेजबान पर DtoD nvvp साथ समय 8.6 एमएस

@talonmies की रिपोर्ट है कि memcpy प्रतियां बाइट द्वारा बाइट जो भार और दुकानों के साथ अक्षम है। मैं अभी भी गणना 3.0 को लक्षित कर रहा हूं इसलिए मैं डिवाइस पर cudaMemcpy का परीक्षण नहीं कर सकता।