-2
共有配列にデータをロードした後、次のコードで、共有メモリ内の配列を合計しようとしています。ロードされた配列のサイズは289です。カーネル内の配列の加算が機能しない
#include <cuda.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<iostream>
#include <stdio.h>
//#include "readmat.cuh"
//#include "mat.h"
#include <device_functions.h>
#include <time.h>
#include <ctime>
//#include "opencv2/highgui/highgui.hpp"
using namespace std;
//using namespace cv;
typedef struct {
size_t X;
size_t Y;
size_t U;
size_t V;
double* elements;
int no_of_elements;
int alpha;
} DataIn;
cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);
__device__ double getelementData(DataIn data, int x, int v, int u, int y)
{
int index = data.X*data.Y*data.U*(v)+data.X*data.Y*(u)+data.X*(y)+x;
return data.elements[index];
}
__global__ void universaladd(DataIn data,int alpha,double* imagedevice)
{
int v = blockIdx.y;
int u = blockIdx.x;
int y = threadIdx.y;
int x = threadIdx.x;
/*
if (x == 0 && y == 0)
{
if (u == 0)
printf(" tooooooooooooooooootal the vale got whenn u=%d v=%d is %f \n", u, v);
if (v == 0)
printf(" tooooooooooooooooootal the vale got whenn u=%d v=%d is %f \n", u, v);
}
*/
double temp;
int local_idx = (blockDim.y*threadIdx.x) + threadIdx.y;
extern __shared__ double matrix[];
int m = alpha - 1;
int Y_shift = y*m;
int X_shift = x*m;
if (v < Y_shift && u < X_shift){
matrix[data.Y*x + y] = getelementData(data, x, (data.V - (Y_shift - v)), (data.U - (X_shift - u)), y);
temp = getelementData(data, x, (data.V - (Y_shift - v)), (data.U - (X_shift - u)), y);
//printf("the vale got when y=%d x=%d u=%d v=%d is %f \n", y, x, u, v, temp);
}
else if (v >= Y_shift && u < X_shift){
matrix[data.Y*x + y] = getelementData(data, x, (v - (Y_shift)), (data.U - (X_shift - u)), y);
temp = getelementData(data, x, (v - (Y_shift)), (data.U - (X_shift - u)), y);
//printf("the vale got when y=%d x=%d u=%d v=%d is %f\n ", y, x, u, v, temp);
}
else if (v < Y_shift && u >= X_shift){
matrix[data.Y*x + y] = getelementData(data, x, (data.V - (Y_shift - v)), (u - (X_shift)), y);
temp = getelementData(data, x, (data.V - (Y_shift - v)), (u - (X_shift)), y);
//printf("the vale got when y=%d x=%d u=%d v=%d is %f \n", y, x, u, v, temp);
}
else if (v >= Y_shift && u >= X_shift){
matrix[data.Y*x + y] = getelementData(data, x, (v - (Y_shift)), (u - (X_shift)), y);
temp = getelementData(data, x, (v - (Y_shift)), (u - (X_shift)), y);
//printf("the vale got when y=%d x=%d u=%d v=%d is %f \n", y, x, u, v, temp);
}
//if((u==0 && v==0) && (x>15 && y>15)) // for testing
//printf("the vale got when y=%d x=%d u=%d v=%d is %f \n", y, x, u, v, temp);
//printf("Maaaaaa the vale got when y=%d x=%d u=%d v=%d is %.1f \n", y, x, u, v, matrix[0]);
__syncthreads;
//++++++++++++++++++++++++++++++++++++++++++++++++++++++
/*
//Section 1 for testing
if (x == 0 && y == 0)
{
for (int m = (data.X*data.Y) - 1; m > 0; m--)
{
matrix[0] += matrix[m];
}
}
*/
/*
if (x == 0 && y == 0)
{
if (u < 5 && v <5)
printf(" tooooooooooooooooootal the vale got whenn u=%d v=%d is %f \n", u, v, matrix[data.Y*x + y]);
}
*/
//+++++++++++++++++++++++++++++++++++++++++++++++++++++++++++
//__syncthreads;
///*
//++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++
//sectiion 2
int different = (data.X*data.Y) - 256;
if (local_idx < different-1)
{
printf("the value when x=%d y=%d u=%d v=%d %d %f to val %f is \n",x,y,u,v, local_idx, matrix[256 + local_idx], matrix[local_idx]);
matrix[local_idx] += matrix[256 + local_idx];
}
__syncthreads;
if (local_idx < 128)
{
matrix[local_idx] += matrix[local_idx + 128];
}
__syncthreads;
if (local_idx < 64)
{
matrix[local_idx] += matrix[local_idx + 64];
}
__syncthreads;
if (local_idx < 32) {
matrix[local_idx] += matrix[local_idx + 32];
matrix[local_idx] += matrix[local_idx + 16];
matrix[local_idx] += matrix[local_idx + 8];
matrix[local_idx] += matrix[local_idx + 4];
matrix[local_idx] += matrix[local_idx + 2];
matrix[local_idx] += matrix[local_idx + 1];
}
__syncthreads;
if (local_idx == 0)
imagedevice[data.V*u + v] = matrix[local_idx];
//+++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++
__syncthreads;
//*/
/*
if (x == 0 && y == 0)
{
if (5>u && 5>v)
printf(" tooooooooooooooooootal the vale got whenn u=%d v=%d is %f \n", u, v, matrix[0]);
}
*/
}
/*
__global__ void intergershift(DataIn data)
{
int v = threadIdx.x;
int u = threadIdx.y;
int y = blockIdx.x;
int Height = blockDim.y;
extern __shared__ double* dataShared[];
double *** dataPoint = (double***)&dataShared;
double *** dataPointShifted = (double***)&dataShared[data.V];
double ** dataElements = (double**)&dataShared[2 * data.V];
dataPoint[v] = &dataShared[2 * data.V + Height*v];
}*/
int main()
{
clock_t begin = clock();
time_t start, end;
int elements;
int numberofdimension;
//const char *file = "Bracelet.mat";
//const size_t* dimepointer;
//readmat thismat(file);
//numberofdimension = thismat.getnumbrofdimensions();
//dimepointer = thismat.dimensionpointer();
size_t X, Y, U, V;
X = 17;
Y = 17;
U = 512;
V = 320;
// Dimensions end
DataIn data;
data.U = U;
data.Y = Y;
data.X = X;
data.V = V;
size_t size = X*Y*U*V*sizeof(double);
data.no_of_elements = X*Y*U*V;
cudaError_t status;
double * dataarray = new double[X*Y*U*V];
for (int k = 0; k < U*V*X*Y; k++)
dataarray[k] = 225;
short * Device_data;
cout << "the size is" << sizeof(Device_data) << endl;
status = cudaSetDevice(0);
status = cudaMalloc((void**)&data.elements, size);
status = cudaMemcpy(data.elements, dataarray, size, cudaMemcpyHostToDevice);
if (status != cudaSuccess) {
fprintf(stderr, "Memory copyind original data failed");
cudaFree(data.elements);
//cudaFree(arrangeddata);
}
/*
for (int t = 0; t < 10; t++)
cout << *(thismat.getarraypointer()+t)<<" ";
cout << endl << "original data printed" << endl;;
*/
double* image = new double[U*V];
double *imagedevice;
status = cudaMalloc((void**)&imagedevice, sizeof(double)*U*V);
//__global__ void universaladd(DataIn data,int alpha,double* imagedevice)
dim3 dimBlock(data.X, data.Y);
dim3 dimGrid(data.U, data.V);
universaladd << <dimGrid, dimBlock, sizeof(double)*X*Y >> >(data, 1, imagedevice);
status = cudaGetLastError();
fprintf(stderr, "Launch status: %s\n", cudaGetErrorString(status));
status = cudaDeviceSynchronize();
fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching\n", status);
status = cudaMemcpy(image, imagedevice, sizeof(double)*U*V, cudaMemcpyDeviceToHost);
if (status != cudaSuccess) {
fprintf(stderr, "Memory copyind original data failed");
cudaFree(data.elements);
}
//cudaFree(arrangeddata);
for (int t = 0; t < 15; t++)
cout << image[t] << " ";
cout << endl << "data printed device";
int k;
std::cin >> k;
return 0;
}
// Helper function for using CUDA to add vectors in parallel.
ただし、ここでは希望の結果が得られません。私はセクション1でループを試してみましたが、結果はすべてのスレッドを実行する前にカーネルが終了しても問題ありません。デバイス同期は私にエラーコード4を与えます。しかし、デバイス同期の前に、カーネルの起動に関連するエラーはありません。
私がセクション2のコードを試しているとき、私は正しく和を得ていません。カーネルの起動やデバイスの同期に関連するエラーメッセージはありません。しかし、その部分が別のコードで試してみるとうまく動作します。ここでロードされるデータは1から255までのピクセル値です。ループが使用されているときに合計が正しいので、ロードされたデータに問題はないと思います。また、ここでは値が255か否定値かをチェックしています。コンソールでは、スレッドは0から255の間にない値をロードします。これは、コンソールに値が表示されないためです。問題は合計を間違って取得するだけでなく、一部の値が異常です。グリッド寸法は320x512（data.Uxdata.V）、ブロック寸法は17x17（data.Xxdata.Y）です。したがって、行列[]の値が有効であれば、最大値は255 * 289でなければなりませんが、負の値と非常に大きな値があります。問題は縮小と思われますが、
ループで起動時に出力が
である（セクション1）セクション2は、デバイスからホストへのメモリコピーが、私は、タスクを完了する前に停止するカーネルに起因すると思われなかった。ここで
Launch status: no error
cudaDeviceSynchronize returned error code 0 after launching
65025 65025 65025 65025 65025 65025 65025 65025 65025 65025 65025 65025 65025 65025 65025
data printed device
をコメントしています。しかし、カーネル内に印刷された値は、私がmatlabに対してチェックしたときに正しいものです。
ここ[UPDATE]増加WDDM TDRの遅延ループが細かい値で実行した後には正しいが、セクション2が依然として所望results.since 1つの1Dブロックで起動されるだけでアレイを有する微細クションsection2ランを与えないと1xsizeスレッド私は、問題は、起動時の設定と縮小との非互換性に関連していると思います。
アレイ低減部1で起動
Launch status: no error
cudaDeviceSynchronize returned error code 0 after launching
85275 7650 27675 58500 27450 103050 30375 17775 18000 12825 24750 95625 15975 68175 7425
コメントされ、ここで
をホストするデバイスからコピー画像に関するエラーはありません[編集]
これはとカーネルでありますメイン。以前はマットファイルからデータをロードしていましたが、データを225値で埋めました。
[MCVE] – talonmies
@talonmiesなしであなたを助けることができる人は誰も私を指摘してくれてありがとう。私は私の質問を更新しました。 –
あなたの編集がどのように役立つか分かりません。誰かがコンパイルして実行できる最短で簡単な例を生成する必要があります。 – talonmies