मेरे पास कक्षा असाइनमेंट के लिए एक साधारण CUDA समस्या थी, लेकिन प्रोफेसर ने साझा स्मृति का उपयोग करके उसी एल्गोरिदम को लागू करने के लिए एक वैकल्पिक कार्य जोड़ा। मैं समय सीमा से पहले इसे पूरा करने में असमर्थ था (जैसा कि में एक सप्ताह पहले था) लेकिन मैं अभी भी उत्सुक हूं इसलिए अब मैं इंटरनेट से पूछूंगा;)।CUDA: एक बड़े-आश 2 डी सरणी पर साझा स्मृति
मूल असाइनमेंट क्रमशः और सीयूडीए दोनों में लाल-काले लगातार ओवर-छूट के एक बेस्टर्ड संस्करण को कार्यान्वित करना था, सुनिश्चित करें कि आपको दोनों में एक ही परिणाम मिला और फिर गति की तुलना करें। जैसे मैंने कहा, साझा स्मृति के साथ ऐसा करना एक वैकल्पिक + 10% ऐड-ऑन था।
मैं अपने कामकाजी संस्करण और छद्म कोड को पोस्ट करने जा रहा हूं जो मैंने करने का प्रयास किया है क्योंकि मेरे पास इस समय मेरे हाथों में कोड नहीं है, लेकिन अगर किसी को किसी को चाहिए तो मैं इसे वास्तविक कोड के साथ अपडेट कर सकता हूं यह।
इससे पहले कि कोई इसे कहता है: हाँ, मुझे पता है कि CUtil का उपयोग लंगड़ा है, लेकिन यह तुलना और टाइमर को आसान बना देता है।
कार्य वैश्विक स्मृति संस्करण:
#define N 1024
__global__ void kernel(int *d_A, int *d_B, int width) {
//assuming width is 64 because that's the biggest number I can make it
//each MP has 48KB of shared mem, which is 12K ints, 32 threads/warp, so max 375 ints/thread?
__shared__ int A_sh[3][66];
//get x and y index and turn it into linear index
for(i=0; i < width+2; i++) //have to load 2 extra values due to the -1 and +1 in algo
A_sh[index_y%3][i] = d_A[index+i-1]; //so A_sh[index_y%3][0] is actually d_A[index-1]
__syncthreads(); //and hope that previous and next row have been loaded by other threads in the block?
//ignore boundary conditions because it's pseudocode
for(i=0; i < width; i++)
d_B[index+i] = A_sh[index_y%3][i] + A_sh[index_y%3][i+2] + A_sh[index_y%3-1][i+1] + A_sh[index_y%3+1][i+1];
}
main(){
//same init as above until threads/grid init
dim3 threadsperblk(32,16);
dim3 numblks(32,64);
kernel<<<numblks,threadsperblk>>>(d_A,d_B,64);
//rest is the same
}
इस साझा मेम कोड दुर्घटनाओं ("लांच के कारण विफल:
#include <stdlib.h>
#include <stdio.h>
#include <cutil_inline.h>
#define N 1024
__global__ void kernel(int *d_A, int *d_B) {
unsigned int index_x = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int index_y = blockIdx.y * blockDim.y + threadIdx.y;
// map the two 2D indices to a single linear, 1D index
unsigned int grid_width = gridDim.x * blockDim.x;
unsigned int index = index_y * grid_width + index_x;
// check for boundaries and write out the result
if((index_x > 0) && (index_y > 0) && (index_x < N-1) && (index_y < N-1))
d_B[index] = (d_A[index-1]+d_A[index+1]+d_A[index+N]+d_A[index-N])/4;
}
main (int argc, char **argv) {
int A[N][N], B[N][N];
int *d_A, *d_B; // These are the copies of A and B on the GPU
int *h_B; // This is a host copy of the output of B from the GPU
int i, j;
int num_bytes = N * N * sizeof(int);
// Input is randomly generated
for(i=0;i<N;i++) {
for(j=0;j<N;j++) {
A[i][j] = rand()/1795831;
//printf("%d\n",A[i][j]);
}
}
cudaEvent_t start_event0, stop_event0;
float elapsed_time0;
CUDA_SAFE_CALL(cudaEventCreate(&start_event0));
CUDA_SAFE_CALL(cudaEventCreate(&stop_event0));
cudaEventRecord(start_event0, 0);
// sequential implementation of main computation
for(i=1;i<N-1;i++) {
for(j=1;j<N-1;j++) {
B[i][j] = (A[i-1][j]+A[i+1][j]+A[i][j-1]+A[i][j+1])/4;
}
}
cudaEventRecord(stop_event0, 0);
cudaEventSynchronize(stop_event0);
CUDA_SAFE_CALL(cudaEventElapsedTime(&elapsed_time0,start_event0, stop_event0));
h_B = (int *)malloc(num_bytes);
memset(h_B, 0, num_bytes);
//ALLOCATE MEMORY FOR GPU COPIES OF A AND B
cudaMalloc((void**)&d_A, num_bytes);
cudaMalloc((void**)&d_B, num_bytes);
cudaMemset(d_A, 0, num_bytes);
cudaMemset(d_B, 0, num_bytes);
//COPY A TO GPU
cudaMemcpy(d_A, A, num_bytes, cudaMemcpyHostToDevice);
// create CUDA event handles for timing purposes
cudaEvent_t start_event, stop_event;
float elapsed_time;
CUDA_SAFE_CALL(cudaEventCreate(&start_event));
CUDA_SAFE_CALL(cudaEventCreate(&stop_event));
cudaEventRecord(start_event, 0);
// TODO: CREATE BLOCKS AND THREADS AND INVOKE GPU KERNEL
dim3 block_size(256,1,1); //values experimentally determined to be fastest
dim3 grid_size;
grid_size.x = N/block_size.x;
grid_size.y = N/block_size.y;
kernel<<<grid_size,block_size>>>(d_A,d_B);
cudaEventRecord(stop_event, 0);
cudaEventSynchronize(stop_event);
CUDA_SAFE_CALL(cudaEventElapsedTime(&elapsed_time,start_event, stop_event));
//COPY B BACK FROM GPU
cudaMemcpy(h_B, d_B, num_bytes, cudaMemcpyDeviceToHost);
// Verify result is correct
CUTBoolean res = cutComparei((int *)B, (int *)h_B, N*N);
printf("Test %s\n",(1 == res)?"Passed":"Failed");
printf("Elapsed Time for Sequential: \t%.2f ms\n", elapsed_time0);
printf("Elapsed Time for CUDA:\t%.2f ms\n", elapsed_time);
printf("CUDA Speedup:\t%.2fx\n",(elapsed_time0/elapsed_time));
cudaFree(d_A);
cudaFree(d_B);
free(h_B);
cutilDeviceReset();
}
साझा स्मृति संस्करण के लिए, यह है कि क्या मैं अब तक की कोशिश की है है अनिर्दिष्ट त्रुटि ") क्योंकि मैंने अभी तक सभी सीमा परिस्थितियों को नहीं पकड़ा है, लेकिन मैं इसके बारे में चिंतित नहीं हूं जितना चीजें पाने के लिए सही तरीका ढूंढना। मुझे लगता है कि मेरा कोड सही पथ (विशेष रूप से एसडीके उदाहरणों की तुलना में) के लिए बहुत जटिल है, लेकिन मैं इसे करने का दूसरा तरीका भी नहीं देख सकता क्योंकि मेरी सरणी साझा किए गए mem में सभी उदाहरणों की तरह फिट नहीं है I मिल सकता है।
और स्पष्ट रूप से, मुझे यकीन है कि यह है कि मेरे हार्डवेयर पर बहुत तेजी से किया जाएगा नहीं कर रहा हूँ (GTX 560 ती - 0.121ms में वैश्विक स्मृति संस्करण चलाता है), लेकिन मैं अपने आप से यह पहली बार साबित करने के लिए की जरूरत है: पी
संपादित करें 2: भविष्य में इस पर चलने वाले किसी भी व्यक्ति के लिए, यदि आप कुछ साझा मेमोरी करना चाहते हैं तो उत्तर में कोड एक अच्छा प्रारंभिक बिंदु है।
इसके बारे में उस तरह से नहीं सोचा था धन्यवाद,: कि कैसे लागू किया जा सकता का एक ठोस उदाहरण के लिए। सवाल यह है कि, मैं ब्लॉक में धागे को एक-दूसरे पर चलने से कैसे रोकूं? मान लें कि मेरे पास ब्लॉक में 2 धागे हैं, और थ्रेड 2 पंक्ति (एफ) लोड करना चाहता है जबकि थ्रेड 1 अभी भी पंक्ति (सी) पर काम कर रहा है? या क्या मुझे कोड को प्रति थ्रेड 1 थ्रेड रखने के लिए कोड बदलना चाहिए और फिर कई ब्लॉक हैं? – a5ehren
@ a5ehren: एक इंट्रा-ब्लॉक सिंक्रनाइज़ेशन आदिम है जिसे __syncthreads() कहा जाता है जिसे आप थ्रेड सिंक्रनाइज़ करने के लिए उपयोग कर सकते हैं। आदर्श रूप में आप प्रति ब्लॉक 32 धागे के कुछ दौर चाहते हैं, और इनपुट स्थान की पंक्ति चौड़ाई को कवर करने के लिए आवश्यक कई ब्लॉक हैं। यदि आप कुछ और सहायता चाहते हैं तो मैं जवाब में थोड़ा छद्म कोड जोड़ सकता हूं। – talonmies
तो क्या प्रत्येक थ्रेड पंक्ति के अपने हिस्से को लोड करेगा, इसे सिंक करेगा, और फिर मान लें कि ऊपर और नीचे पंक्तियों पर काम करने वाले धागे हैं? मुझे लगता है कि कुछ छद्म कोड मदद करेंगे: पी – a5ehren