質問をすることでしか得られない、回答やアドバイスがある。

15分調べてもわからないことは、質問しよう!

新規登録して質問してみよう
ただいま回答率
85.50%
CUDA

CUDAは並列計算プラットフォームであり、Nvidia GPU(Graphics Processing Units)向けのプログラミングモデルです。CUDAは様々なプログラミング言語、ライブラリ、APIを通してNvidiaにインターフェイスを提供します。

C++

C++はC言語をもとにしてつくられた最もよく使われるマルチパラダイムプログラミング言語の1つです。オブジェクト指向、ジェネリック、命令型など広く対応しており、多目的に使用されています。

Q&A

解決済

1回答

3894閲覧

CUDAでのクラス利用時の他オブジェクトからのデバイスメモリを参照できない

退会済みユーザー

退会済みユーザー

総合スコア0

CUDA

CUDAは並列計算プラットフォームであり、Nvidia GPU(Graphics Processing Units)向けのプログラミングモデルです。CUDAは様々なプログラミング言語、ライブラリ、APIを通してNvidiaにインターフェイスを提供します。

C++

C++はC言語をもとにしてつくられた最もよく使われるマルチパラダイムプログラミング言語の1つです。オブジェクト指向、ジェネリック、命令型など広く対応しており、多目的に使用されています。

0グッド

0クリップ

投稿2019/04/27 04:54

発生している問題

下記のコードはCUDAでクラスを用いて行列の内積を行うコードです。
予想どおり動くならば結果は

[[20.000000,23.000000,26.000000,29.000000]
[56.000000,68.000000,80.000000,92.000000]]

となるはずです。

しかし結果は

[[0.000000,0.000000,0.000000,0.000000]
[0.000000,0.000000,0.000000,0.000000]]

となってしまい計算できていません。

値が全て0なのは領域確保後ゼロクリアしているのでおかしくはないのですが
カーネル関数のDot_Kernel内のデバイス側メモリのポインタmat1とmat2が
参照できていない事までは突き止めました。

しなぜ参照できないのか?
どうすれば参照できるのか?

は私の知識だけではどうしても解決できないので
どうかご教授よろしくお願いします。
また、私はまだCUDAとC++を初めて1年も経っていないド素人ですので
めちゃくちゃなこと言ってたらすいません。

###コード
本コードは3つのファイルから構成されています。
また、それらのファイルは全て同じディレクトリに存在します。

main.cpp

c++

1///////////////////////////////////////////////////////////////////////////////// 2// 3// 本ファイルではcudaを用いた行列クラス同士の内積処理のメインルーチンを記す。 4// 5///////////////////////////////////////////////////////////////////////////////// 6 7#include "matrix.cuh" 8 9int main(){ 10 11 Matrix2d* mat1 = new Matrix2d(2,3);//2行3列の行列オブジェクト作成(入力1) 12 Matrix2d* mat2 = new Matrix2d(3,4);//3行4列の行列オブジェクト作成(入力2) 13 Matrix2d* out = new Matrix2d(2,4);//2行4列の行列オブジェクト作成(内積結果保存用) 14 15 mat1->RangeInt();//0から要素数までの値挿入 16 ////////////////////////////////////// 17 // 現時点でのmat1の中身 18 // [[0.000000,1.000000,2.000000] 19 // [3.000000,4.000000,5.000000]] 20 ////////////////////////////////////// 21 22 mat2->RangeInt(); 23 /////////////////////////////////////////////////// 24 // 現時点でのmat2の中身 25 // [[0.000000,1.000000,2.000000,3.000000] 26 // [4.000000,5.000000,6.000000,7.000000] 27 // [8.000000,9.000000,10.000000,11.000000]] 28 /////////////////////////////////////////////////// 29 30 out->Dot(mat1,mat2);// out = mat1 dot mat2 (内積計算) 31 32 out->Print();// オブジェクトoutの中身表示 33 ///////////////////////////////////////////////// 34 // ▲問題発生▲ 35 // 36 // 内積結果 37 // [[0.000000,0.000000,0.000000,0.000000] 38 // [0.000000,0.000000,0.000000,0.000000]] 39 // 40 // 内積計算の結果が入っていない 41 // 42 ///////////////////////////////////////////////// 43 44 delete mat1;//オブジェクト作成で確保した領域開放 45 delete mat2;//オブジェクト作成で確保した領域開放 46 delete out;//オブジェクト作成で確保した領域開放 47 48 return 0; 49}

matrix.cuh

CUDA

1////////////////////////////////////////////////////// 2// 3// 本ファイルでは二次元行列クラスの定義を記す。 4// 5////////////////////////////////////////////////////// 6#ifndef MATRIX_CUH 7#define MATRIX_CUH 8 9#include <stdio.h> 10#include <stdlib.h> 11 12//1ワープあたりのスレッド数 13const int WORP = 32; 14 15class Matrix2d; 16 17//二次元行列クラス 18class Matrix2d{ 19 private: 20 //見かけ上の行数 21 int ColTrue; 22 //見かけ上の列数 23 int RowTrue; 24 25 //カーネル実行時のgrids/blocksとblocks/threadsのつじつまを合わせるための行数 26 int ColAll; 27 //カーネル実行時のgrids/blocksとblocks/threadsのつじつまを合わせるための列数 28 int RowAll; 29 30 //デバイスメモリのポインタを保存するためのメンバ変数 31 double* Element; 32 33 public: 34 //コンストラクタ 35 Matrix2d(int Col,int Row); 36 //デストラクタ 37 ~Matrix2d(); 38 39 //オブジェクト内の行列表示 40 void Print()const; 41 //0~要素数までの数値で埋める 42 void RangeInt(); 43 44 //見かけ上の行数の行数のゲッター 45 int GetColTrue()const; 46 //見かけ上の行数の列数のゲッター 47 int GetRowTrue()const; 48 //カーネル実行時のgrids/blocksとblocks/threadsのつじつまを合わせるための行数のゲッター 49 int GetColAll()const; 50 //カーネル実行時のgrids/blocksとblocks/threadsのつじつまを合わせるための列数のゲッター 51 int GetRowAll()const; 52 //デバイスメモリのポインタを保存するためのメンバ変数のゲッター 53 double* GetElement()const; 54 55 //内積処理 56 void Dot(const Matrix2d* mat1,const Matrix2d* mat2); 57}; 58 59#endif

matrix.cu

CUDA

1//////////////////////////////////////////////////////////////////////////////////////////// 2// 3// 本ファイルではmatrix.cuhファイルで定義したクラスの実装とカーネル関数の実装を記す。 4// 5//////////////////////////////////////////////////////////////////////////////////////////// 6 7#include "matrix.cuh" 8 9//////////////////////////// 10// 11// ここからカーネル関数 12// 13//////////////////////////// 14 15//デバイス側の二次元行列の中身表示 16__global__ 17void Print_Kernel(int Col,int Row,double* Element){ 18 printf("["); 19 for(int i=0;i<Col;i++){ 20 if(i != 0)printf(" "); 21 printf("["); 22 for(int j=0;j<Row;j++){ 23 printf("%lf",Element[j+Row*i]); 24 if(j != Row-1)printf(","); 25 } 26 if(i != Col-1){ 27 printf("]\n"); 28 }else{ 29 printf("]"); 30 } 31 } 32 printf("]\n"); 33} 34 35//内積処理 36__global__ 37void Dot_Kernrl(double* out,const double* mat1,const double* mat2,int mat1_row,int mat2_row,int out_row){ 38 int x = blockIdx.x * blockDim.x + threadIdx.x; 39 int y = blockIdx.y * blockDim.y + threadIdx.y; 40 41 double sum = 0; 42 for(int i=0;i<mat1_row;i++){ 43 sum += mat1[i+(mat1_row)*y] * mat2[x+(mat2_row)*i];// ←どうやらここでmat1とmat2が参照できていない 44 } 45 out[x+(out_row)*y] = sum;//←ここをout[x+(out_row)*y] = 123とするとちゃんとoutの全てに123が入るのでoutへのアクセスは問題ない 46 47 /////////////////////////////////////////// 48 // mat1とmat2の中身が参照できていない。 49 /////////////////////////////////////////// 50} 51 52////////////////////////////// 53// 54// ここからMatrix2dクラス 55// 56////////////////////////////// 57//コンストラクタ 58Matrix2d::Matrix2d(int Col,int Row){ 59 this->ColTrue = Col; 60 this->RowTrue = Row; 61 62 this->ColAll = WORP*((this->ColTrue/WORP)+1); 63 this->RowAll = WORP*((this->RowTrue/WORP)+1); 64 65 cudaMalloc((double**)&this->Element,sizeof(double)*this->ColAll*this->RowAll); 66 67 double* zeros; 68 zeros = (double*)malloc(sizeof(double)*this->ColAll*this->RowAll); 69 for(int i=0;i<this->ColAll*this->RowAll;i++){ 70 zeros[i] = 0; 71 } 72 73 cudaMemcpy(this->Element,zeros,sizeof(double)*this->ColAll*this->RowAll,cudaMemcpyHostToDevice); 74 free(zeros); 75 zeros = NULL; 76} 77 78//デストラクタ 79Matrix2d::~Matrix2d(){ 80 cudaFree(this->Element); 81 this->Element = NULL; 82} 83 84//オブジェクト内の行列表示 85void Matrix2d::Print()const{ 86 Print_Kernel<<<1,1>>>(this->ColTrue,this->RowTrue,this->Element); 87 cudaDeviceSynchronize(); 88} 89 90//0~要素数までの数値で埋める 91void Matrix2d::RangeInt(){ 92 double* ran = (double*)malloc(sizeof(double)*this->ColTrue*this->RowTrue); 93 for(int i=0;i<this->ColTrue*this->RowTrue;i++){ 94 ran[i] = i; 95 } 96 cudaMemcpy(this->Element,ran,sizeof(double)*this->ColTrue*this->RowTrue,cudaMemcpyHostToDevice); 97 free(ran); 98} 99 100//見かけ上の行数の行数のゲッター 101int Matrix2d::GetColTrue()const{ 102 return this->ColTrue; 103} 104 105//見かけ上の行数の列数のゲッター 106int Matrix2d::GetRowTrue()const{ 107 return this->RowTrue; 108} 109 110//カーネル実行時のgrids/blocksとblocks/threadsのつじつまを合わせるための行数のゲッター 111int Matrix2d::GetColAll()const{ 112 return this->ColAll; 113} 114 115//カーネル実行時のgrids/blocksとblocks/threadsのつじつまを合わせるための列数のゲッター 116int Matrix2d::GetRowAll()const{ 117 return this->RowAll; 118} 119 120//デバイスメモリのポインタを保存するためのメンバ変数のゲッター 121double* Matrix2d::GetElement()const{ 122 return this->Element; 123} 124 125//内積処理 126void Matrix2d::Dot(const Matrix2d* mat1,const Matrix2d* mat2){ 127 dim3 grids; 128 dim3 blocks; 129 130 grids.x = this->RowAll/WORP; 131 grids.y = this->ColAll/WORP; 132 grids.z = 1; 133 134 blocks.x = WORP; 135 blocks.y = WORP; 136 blocks.z = 1; 137 138 Dot_Kernrl<<<grids,blocks>>>(this->Element,mat1->GetElement(),mat2->GetElement(),mat1->GetRowAll(),mat2->GetRowAll(),this->RowAll); 139 cudaDeviceSynchronize(); 140}

コンパイル方法

terminal

1nvcc main.cpp matrix.cu

上記コマンドを実行後このディレクトリ内にa.outという実行ファイルが作成される。

補足情報(FW/ツールのバージョンなど)

os:linux mint18.3 cinnamon
cuda version:9.0
GPU:nVidia GeForce GTX1060 3Gb

気になる質問をクリップする

クリップした質問は、後からいつでもMYページで確認できます。

またクリップした質問に回答があった際、通知やメールを受け取ることができます。

バッドをするには、ログインかつ

こちらの条件を満たす必要があります。

guest

回答1

0

ベストアンサー

どうやら要素番号の計算方法で勘違いしていただけのようです。
アクセスはできていました。
とんだ勘違いでお騒がせしてしまい失礼いたしました。

修正後のmatrix.cu

CUDA

1//////////////////////////////////////////////////////////////////////////////////////////// 2// 3// 本ファイルではmatrix.cuhファイルで定義したクラスの実装とカーネル関数の実装を記す。 4// 5//////////////////////////////////////////////////////////////////////////////////////////// 6 7#include "matrix.cuh" 8 9//////////////////////////// 10// 11// ここからカーネル関数 12// 13//////////////////////////// 14 15//デバイス側の二次元行列の中身表示 16__global__ 17void Print_Kernel(int Col,int Row,int RowAll,double* Element){ 18 printf("["); 19 for(int i=0;i<Col;i++){ 20 if(i != 0)printf(" "); 21 printf("["); 22 for(int j=0;j<Row;j++){ 23 printf("%lf",Element[j+RowAll*i]);// <- 変更点 24 if(j != Row-1)printf(","); 25 } 26 if(i != Col-1){ 27 printf("]\n"); 28 }else{ 29 printf("]"); 30 } 31 } 32 printf("]\n"); 33} 34 35//内積処理 36__global__ 37void Dot_Kernrl(double* out,const double* mat1,const double* mat2,int mat1_row,int mat2_row,int out_row){ 38 int x = blockIdx.x * blockDim.x + threadIdx.x; 39 int y = blockIdx.y * blockDim.y + threadIdx.y; 40 41 double sum = 0; 42 for(int i=0;i<mat1_row;i++){ 43 sum += mat1[i+(mat1_row)*y] * mat2[x+(mat2_row)*i];// ←どうやらここでmat1とmat2が参照できていない 44 } 45 out[x+(out_row)*y] = sum;//←ここをout[x+(out_row)*y] = 123とするとちゃんとoutの全てに123が入るのでoutへのアクセスは問題ない 46 47 /////////////////////////////////////////// 48 // mat1とmat2の中身が参照できていない。 49 /////////////////////////////////////////// 50} 51 52////////////////////////////// 53// 54// ここからMatrix2dクラス 55// 56////////////////////////////// 57//コンストラクタ 58Matrix2d::Matrix2d(int Col,int Row){ 59 this->ColTrue = Col; 60 this->RowTrue = Row; 61 62 this->ColAll = WORP*((this->ColTrue/WORP)+1); 63 this->RowAll = WORP*((this->RowTrue/WORP)+1); 64 65 cudaMalloc((double**)&this->Element,sizeof(double)*this->ColAll*this->RowAll); 66 67 double* zeros; 68 zeros = (double*)malloc(sizeof(double)*this->ColAll*this->RowAll); 69 for(int i=0;i<this->ColAll*this->RowAll;i++){ 70 zeros[i] = 0; 71 } 72 73 cudaMemcpy(this->Element,zeros,sizeof(double)*this->ColAll*this->RowAll,cudaMemcpyHostToDevice); 74 free(zeros); 75 zeros = NULL; 76} 77 78//デストラクタ 79Matrix2d::~Matrix2d(){ 80 cudaFree(this->Element); 81 this->Element = NULL; 82} 83 84//オブジェクト内の行列表示 85void Matrix2d::Print()const{ 86 Print_Kernel<<<1,1>>>(this->ColTrue,this->RowTrue,this->RowAll,this->Element); 87 cudaDeviceSynchronize(); 88} 89 90//0~要素数までの数値で埋める 91void Matrix2d::RangeInt(){ 92 double* ran = (double*)malloc(sizeof(double)*this->ColAll*this->RowAll); // <- 変更点 93 int count=0;// <- 変更点 94 for(int i=0;i<this->ColTrue;i++){// <- 変更点 95 for(int j=0;j<this->RowTrue;j++){// <- 変更点 96 ran[j+this->RowAll*i] = count;// <- 変更点 97 count++;// <- 変更点 98 } 99 } 100 cudaMemcpy(this->Element,ran,sizeof(double)*this->ColAll*this->RowAll,cudaMemcpyHostToDevice);// <- 変更点 101 free(ran); 102} 103 104//見かけ上の行数の行数のゲッター 105int Matrix2d::GetColTrue()const{ 106 return this->ColTrue; 107} 108 109//見かけ上の行数の列数のゲッター 110int Matrix2d::GetRowTrue()const{ 111 return this->RowTrue; 112} 113 114//カーネル実行時のgrids/blocksとblocks/threadsのつじつまを合わせるための行数のゲッター 115int Matrix2d::GetColAll()const{ 116 return this->ColAll; 117} 118 119//カーネル実行時のgrids/blocksとblocks/threadsのつじつまを合わせるための列数のゲッター 120int Matrix2d::GetRowAll()const{ 121 return this->RowAll; 122} 123 124//デバイスメモリのポインタを保存するためのメンバ変数のゲッター 125double* Matrix2d::GetElement()const{ 126 return this->Element; 127} 128 129//内積処理 130void Matrix2d::Dot(const Matrix2d* mat1,const Matrix2d* mat2){ 131 dim3 grids; 132 dim3 blocks; 133 134 grids.x = this->RowAll/WORP; 135 grids.y = this->ColAll/WORP; 136 grids.z = 1; 137 138 blocks.x = WORP; 139 blocks.y = WORP; 140 blocks.z = 1; 141 142 Dot_Kernrl<<<grids,blocks>>>(this->Element,mat1->GetElement(),mat2->GetElement(),mat1->GetRowAll(),mat2->GetRowAll(),this->RowAll); 143 cudaDeviceSynchronize(); 144}

投稿2019/04/27 07:07

退会済みユーザー

退会済みユーザー

総合スコア0

バッドをするには、ログインかつ

こちらの条件を満たす必要があります。

あなたの回答

tips

太字

斜体

打ち消し線

見出し

引用テキストの挿入

コードの挿入

リンクの挿入

リストの挿入

番号リストの挿入

表の挿入

水平線の挿入

プレビュー

15分調べてもわからないことは
teratailで質問しよう!

ただいまの回答率
85.50%

質問をまとめることで
思考を整理して素早く解決

テンプレート機能で
簡単に質問をまとめる

質問する

関連した質問