Как я уже упоминал в комментариях, ваше понимание указателей неверно. Я внес изменения во многих местах в вашем коде, чтобы решить эту проблему. Я отметил большинство из них комментарием // mod, но я, возможно, пропустил некоторые.
Кроме того, ваше ядро просто не может отслеживать элементы, когда несколько потоков могут обновлять одно и то же местоположение. Один из способов разобраться в этом - использовать атомарность (которую я продемонстрировал). Существуют различные другие подходы, такие как параллельное сокращение, но ни один из них не является тривиальным изменением ядра. Кроме того, логика вашего ядра была нарушена несколькими способами.
Далее следует наименьшее количество модификаций, которые я мог бы внести в ваш код, чтобы получить что-то разумное. Существует несколько параметров компиляции, которые вы можете использовать для изучения поведения ядра:
- без переключателя - близко к ядру, но оно не будет работать правильно
-DUSE_ATOMICS продемонстрирует модификацию вашего ядра для правильного подсчета. -DUSE_ALT_KERNEL исследует другой подход к логике ядра: назначьте один поток на каждый блок гистограммы и проследите, чтобы каждый поток проходил по всему массиву, отслеживая элементы, которые принадлежат этому бину. Поскольку только один поток записывает в каждый результат bin, нет необходимости в атомарности. Однако у нас может быть только столько потоков (с этой тривиальной реализацией), сколько существует бинов. Без особых трудностей этот метод, вероятно, можно было бы расширить до одного деформирования на бин, с использованием деформации основы для окончательного уменьшения уровня деформации , прежде чем один поток записал окончательные результаты в бин. Это несколько повысит эффективность доступа к памяти. Однако это также привнесет в ядро сложность, которую вы, вероятно, еще не изучили.
Вот код:
$ cat t316.cu
#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#define MAXR 16
#define MAXC 16
#define BINS 32
#define N (MAXR*MAXC)
__global__ void count(int *arrayONE_d, int *occurrences_d, int *occurrences_final_d) {
//provide unique thread ID
int idx = threadIdx.x + blockIdx.x * blockDim.x;
#ifndef USE_ALT_KERNEL
if(idx < N) {
//for(k=0; k < MAXR*MAXC; k++) {
for(int j=0; j<32; j++) {
if(arrayONE_d[idx]==occurrences_d[j]){
#ifndef USE_ATOMICS
occurrences_final_d[j]++;
#else
atomicAdd(occurrences_final_d+j, 1);
#endif
}
else {}
}
}
#else
// use one thread per histo bin
if (idx < BINS){
int count = 0;
int myval = occurrences_d[idx];
for (int i = 0; i < N; i++) if (arrayONE_d[i] == myval) count++;
occurrences_final_d[idx] = count;
}
#endif
}
int main(void) {
//const int N = MAXR*MAXC;
int arr1_h[MAXR][MAXC];
//int *occurrences_h[0][32];
//creating arrays for the device (GPU)
//int *arr1_d;
int occurrences_h[32]; // mod
int *occurrences_d;
int occurrences_final_h[32] = {0}; // mod
int *occurrences_final_d;
int arrayONE_h[256] = {0}; // mod
int *arrayONE_d;
int i, j;
// allocating memory for the arrays on the device
cudaMalloc( (void**) &arrayONE_d, MAXR*MAXC*sizeof(int)); // change to 16384 when using 128x128
cudaMalloc( (void**) &occurrences_d, 32* sizeof(int));
cudaMalloc( (void**) &occurrences_final_d, 32*sizeof(int));
/*
for(i=0; i < 32; i++) {
occurrences_h[i] = i;
}
*/
//Reading in matrix from .txt file and storing it in arr1 on the host (CPU)
// FILE *fp;
// fp =fopen("arrays16.txt","r");
// this loop takes the information from .txt file and puts it into arr1 matrix
for(i=0;i<MAXR;i++) {
for(j=0;j<MAXC;j++)
{
// fscanf(fp,"%d\t", &arr1_h[i][j]);
arr1_h[i][j] = j; // mod
}
}
for(i=0;i<MAXR;i++) {
for(j=0;j<MAXC;j++) {
//printf("d\t", arr1_h[i][j]);
}
}
int x,y;
int z=0;
// this loop flattens the 2d array and makes it a 1d array of length MAXR*MAXC
for(x=0;x<MAXR;x++)
{
for(y=0;y<MAXC;y++)
{
// printf("**%d ",arr1_h[x][y]);
arrayONE_h[z]= arr1_h[x][y]; // mod
z++;
}
}
for(x=0; x < 256; x++) {
// printf("%d\n", arrayONE_h[x]); // mod
//return 0;
}
int length = sizeof(arrayONE_h)/sizeof(arrayONE_h[0]);
printf("**LENGTH = %d\n", length);
// copying the arrays/memory from the host to the device (GPU)
cudaMemcpy(arrayONE_d, arrayONE_h, MAXR*MAXC*sizeof(int), cudaMemcpyHostToDevice); //mod
cudaMemcpy(occurrences_d, occurrences_h, 32*sizeof(int), cudaMemcpyHostToDevice); // mod
cudaMemcpy(occurrences_final_d, occurrences_final_h, 32*sizeof(int), cudaMemcpyHostToDevice); // mod
// how many blocks we will allocate
//dim3 DimGrid();
//how many threads per block we will allocate
#ifndef USE_ALT_KERNEL
dim3 DimBlock(N);
#else
dim3 DimBlock(BINS);
#endif
//kernel launch against the GPU
count<<<1, DimBlock>>>(arrayONE_d,occurrences_d,occurrences_final_d);
//copy the arrays post-computation from the device back to the host (CPU)
cudaMemcpy(occurrences_final_h, occurrences_final_d, 32*sizeof(int), cudaMemcpyDeviceToHost); // mod
cudaMemcpy(occurrences_h, occurrences_d, 32*sizeof(int), cudaMemcpyDeviceToHost); // mod
// some error checking - run this with cuda-memcheck when executing your code
cudaError_t errSync = cudaGetLastError();
cudaError_t errAsync = cudaDeviceSynchronize();
if (errSync != cudaSuccess)
printf("Sync kernel error: %s\n", cudaGetErrorString(errSync));
if (errAsync != cudaSuccess)
printf("Async kernel error: %s\n", cudaGetErrorString(errAsync));
//free up the memory of the device arrays
cudaFree(arrayONE_d);
cudaFree(occurrences_d);
cudaFree(occurrences_final_d);
//print out the number of occurrences of each 0-31 value
for(i=0;i<32;i++) {
printf("%d ",occurrences_final_h[i]);
}
printf("\n");
}
$ nvcc -o t316 t316.cu
$ cuda-memcheck ./t316
========= CUDA-MEMCHECK
**LENGTH = 256
1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1
========= ERROR SUMMARY: 0 errors
$ nvcc -o t316 t316.cu -DUSE_ATOMICS
$ ./t316
**LENGTH = 256
16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16
$ nvcc -o t316 t316.cu -DUSE_ALT_KERNEL
$ cuda-memcheck ./t316
========= CUDA-MEMCHECK
**LENGTH = 256
16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16
========= ERROR SUMMARY: 0 errors
$
В приведенном выше выводе мы видим, что базовое ядро дает неверные результаты. Атомное ядро и альтернативное ядро дают правильные результаты
(Ваш код был изменен для использования синтезированных данных, поэтому ему не нужно открывать файл.)