opencl
परमाणु संचालन
खोज…
वाक्य - विन्यास
int atomic_add (वाष्पशील __global int * p, int val)
अहस्ताक्षरित int atomic_add (अस्थिर __global अहस्ताक्षरित int * p, अहस्ताक्षरित int val)
int atomic_add (वाष्पशील __local int * p, int val)
अहस्ताक्षरित int atomic_add (अस्थिर __local अहस्ताक्षरित int * p, अहस्ताक्षरित int val)
पैरामीटर
पी | वैल |
---|---|
सेल को सूचक | सेल में जोड़ा गया |
टिप्पणियों
प्रदर्शन परमाणु संचालन संख्या और मेमोरी स्पेस पर निर्भर करता है। सीरियल काम करना लगभग हमेशा कर्नेल निष्पादन को धीमा करता है क्योंकि जीपीयू एक सिमड सरणी है और सरणी में प्रत्येक इकाई अन्य इकाइयों की प्रतीक्षा करती है यदि वे एक ही प्रकार का काम नहीं करते हैं।
परमाणु जोड़ें समारोह
int fakeMalloc(__local int * addrCounter,int size)
{
// lock addrCounter
// adds size to addrCounter's pointed cell value
// unlock
// return old value of addrCounter's pointed cell
// serial between all threads visiting -> slow
return atomic_add(addrCounter,size);
}
__kernel void vecAdd(__global float* results )
{
int id = get_global_id(0);
int lid=get_local_id(0);
__local float stack[1024];
__local int ctr;
if(lid==0)
ctr=0;
barrier(CLK_LOCAL_MEM_FENCE);
stack[lid]=0.0f; // parallel operation
barrier(CLK_LOCAL_MEM_FENCE);
int fakePointer=fakeMalloc(&ctr,1); // serial operation
barrier(CLK_LOCAL_MEM_FENCE);
stack[fakePointer]=lid; // parallel operation
barrier(CLK_GLOBAL_MEM_FENCE);
results[id]=stack[lid];
}
पहले तत्वों का उत्पादन:
कभी कभी
192 193 194 195 196 197 198
कभी कभी
0 1 2 3 4 5 6
कभी कभी
128 129 130 131 132 133 134
स्थानीय सीमा = 256 के साथ एक सेटिंग के लिए।
जो भी धागा फर्जी मैल्कॉक का दौरा करता है, वह पहले स्थानीय सेल आईडी को पहले परिणाम सेल में डालता है। उदाहरण के gpu की आंतरिक SIMD और वेवफ्रंट संरचना पड़ोसी परिणाम कोशिकाओं के लिए अपने मूल्यों को रखने के लिए 64 थ्रेड को अनुमति देता है। अन्य उपकरण उस उपकरण के opencl कार्यान्वयन के आधार पर अधिक यादृच्छिक रूप से या पूरी तरह से मान डाल सकते हैं।