Программирование общей памяти CUDA не работает

Question

все:

Я узнаю, как разделяемая память ускоряет процесс программирования графического процессора. Я использую коды ниже, чтобы вычислить квадрат значения каждого элемента плюс квадрат значения среднего его левого и правого соседей. Код работает, однако результат не так, как ожидалось.

Первые 10 результатов распечатываются 0,1,2,3,4,5,6,7,8,9, в то время как я ожидаю результат как 25,2,8, 18,32,50, 72,98,128,162;

Код выглядит следующим образом, со ссылкой на here;

Не могли бы вы рассказать мне, какая часть идет не так? Ваша помощь очень ценится.

#include <stdio.h> 
#include <stdlib.h> 
#include <iostream> 
#include <cuda.h> 

const int N=1024; 

__global__ void compute_it(float *data) 
{ 
int tid = threadIdx.x; 
__shared__ float myblock[N]; 
float tmp; 

// load the thread's data element into shared memory 
myblock[tid] = data[tid]; 

// ensure that all threads have loaded their values into 
// shared memory; otherwise, one thread might be computing 
// on unitialized data. 
__syncthreads(); 

// compute the average of this thread's left and right neighbors 
tmp = (myblock[tid>0?tid-1:(N-1)] + myblock[tid<(N-1)?tid+1:0]) * 0.5f; 
// square the previousr result and add my value, squared 
tmp = tmp*tmp + myblock[tid]*myblock[tid]; 

// write the result back to global memory 
data[tid] = myblock[tid]; 
__syncthreads(); 
    } 

int main(){ 

char key; 

float *a; 
float *dev_a; 

a = (float*)malloc(N*sizeof(float)); 
cudaMalloc((void**)&dev_a,N*sizeof(float)); 

for (int i=0; i<N; i++){ 
a [i] = i; 
} 


cudaMemcpy(dev_a, a, N*sizeof(float), cudaMemcpyHostToDevice); 

compute_it<<<N,1>>>(dev_a); 

cudaMemcpy(a, dev_a, N*sizeof(float), cudaMemcpyDeviceToHost); 


for (int i=0; i<10; i++){ 
std::cout<<a [i]<<","; 
} 

std::cin>>key; 

free (a); 
free (dev_a); 

Answer 1

Одной из самых насущных проблем в коде ядра заключается в следующем:

data[tid] = myblock[tid]; 

Я думаю, что вы, вероятно, имел в виду это:

data[tid] = tmp; 

Кроме того, вы запускаете 1024 блоков одного потока. Это не является особенно эффективным способом использования графического процессора, и это означает, что переменная tid в каждом потоковом блоке равна 0 (и только 0, так как в каждом блоке потоков имеется только один поток.) ​​

С этим подходом существует множество проблем , но одна непосредственная проблема будет встречаться здесь:

tmp = (myblock[tid>0?tid-1:(N-1)] + myblock[tid<31?tid+1:0]) * 0.5f; 

Поскольку tid всегда равен нулю, и, следовательно, никакие другие значения в вашей общей матрице памяти (myblock) не заполнит список, логика в этой линии не может быть разумным. Когда tid равно нулю, вы выбираете myblock[N-1] для первого слагаемого в задании tmp, но myblock[1023] никогда не заселяется ничем.

кажется, что вы не понимаете различные CUDA иерархий:

• сетки является все темы, связанные с запуском ядра
• сетка состоит из threadblocks
• каждый threadblock представляет собой группу нитей, работающих вместе на одном SM
• совместно используемый ресурс памяти является ресурсом за СМ, а не устройство в масштабе ресурс
• __synchthreads() также работает на основе threadblock (не в масштабе устройства)
• threadIdx.x - это встроенная переменная, которая предоставляет уникальный идентификатор потока для всех потоков в потоковом блоке, но не глобально по всей сетке.

Вместо этого вы должны разбить свою проблему на группы блоков замедленного размера (то есть более одного потока). Затем каждый блок будет иметь возможность вести себя так, как это было описано выше. Затем вам понадобится особый случай поведения в начальной и конечной точках (в ваших данных) каждого потокаблока.

Вы также не выполняете должным образом cuda error checking, который рекомендуется, особенно в любое время, когда у вас возникают проблемы с кодом CUDA.

Если вы сделаете изменения я указанный первым в коде ядра, а также изменить порядок ваших блоков и сетевых параметров запуска ядра:

compute_it<<<1,N>>>(dev_a); 

Как указано Kristof, вы получите то, что приближается к что вы хотите, я думаю. Однако вы не сможете удобно масштабировать его за пределы N = 1024 без каких-либо изменений в коде.

Эта строка кода также не правильно:

free (dev_a); 

С dev_a было выделено на устройстве с помощью cudaMalloc вы должны освободить его, как это:

cudaFree (dev_a); 

Answer 2

Поскольку у вас есть только один поток для каждого блок, ваш tid всегда будет 0.

Попробуйте запустить ядро ​​следующим образом: compute_it < < < 1, N >>> (dev_a);

вместо compute_it < < >> (dev_a);