1
512スレッドの線形ブロックでカーネルを起動しています。各スレッドには、合計512×6 * 8 = 24576バイトの6つの倍精度値（2つの3要素ベクトル）が共有メモリに格納されます。私は次のようにすべてのベクトルをラインアップするために、共有の中間の要素へのポインタを作成したいと思います：CUDA：中間共有メモリの場所へのポインタ予期しない動作
__global__ void my_kernel(double *global_data) {
extern __shared__ double shr[];
id = threadIdx.x;
double *X = &shr[id*3];
double *Y = &shr[(id+1)*3];
// Some arithmetic to set X[0:3] ad Y[0:3]
// Now I have a small for loop to compute something for each thread
for (int i = 0; i < 3; i++) {
for (int j=0; j < 3; j++) {
// Some computations involving the X and Y vectors
}
}
私の問題は、ループインデックスを使用して、XとYの値にアクセスしています。最初のループ反復中に以下の動作を説明することができません：
(cuda-gdb) cuda thread
thread (0,0,0)
(cuda-gdb) p shr[0]
$1 = 0.62293193093894383
(cuda-gdb) p &shr[0]
$2 = (@shared double *) 0x0
(cuda-gdb) p X[0]
$3 = 0.62293193093894383
(cuda-gdb) p &X[0]
$4 = (@generic double *) 0x1000000
(cuda-gdb) p X
$5 = (@generic double * @register) 0x1000000
これは正常です。しかし、：
i == 0のときにX [0]にアクセスできますが、X [i]にアクセスできないのはなぜですか？
編集：ここに私の問題を実証する完全な作業例です：
import pycuda.gpuarray as gpuarray
import pycuda.driver as cuda
import pycuda.autoinit
import numpy as np
from pycuda.compiler import SourceModule
from math import pi
mydat = np.arange(12).astype(np.float64)
mydat_gpu = gpuarray.to_gpu(mydat)
mod = SourceModule("""
__global__ void my_kernel(double *mydat) {
extern __shared__ double shr[];
int id = threadIdx.x;
double *X = &shr[(id * 6)];
double *Y = &shr[(id * 6) + 3];
X[0] = mydat[0];
X[1] = mydat[1];
X[2] = mydat[2];
Y[0] = mydat[3];
Y[1] = mydat[4];
Y[2] = mydat[5];
__syncthreads();
double result;
for (int i = 0; i < 3; i++) {
result += X[i] + Y[i];
}
}
""")
my_kernel = mod.get_function("my_kernel")
blk = (1,1,1)
grd = (1,1,1)
my_kernel(mydat_gpu, grid=grd, block=blk, shared=(8*6))
私はデバッグセッションを起動し、この時点で：
cuda-gdb --args python -m pycuda.debug minimal_working_example.py
(cuda-gdb) b my_kernel
Function "my_kernel" not defined.
Make breakpoint pending on future shared library load? (y or [n]) y
Breakpoint 1 (my_kernel) pending.
(cuda-gdb) run
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]
Breakpoint 1, my_kernel(double * @generic)<<<(1,1,1),(1,1,1)>>> (mydat=0x13034a0000)
at kernel.cu:5
5 int id = threadIdx.x;
(cuda-gdb) n
7 double *X = &shr[(id * 6)];
(cuda-gdb) p id
$1 = 0
(cuda-gdb) p id * 6
$2 = 0
(cuda-gdb) n
8 double *Y = &shr[(id * 6) + 3];
(cuda-gdb) p (id * 6) + 3
$3 = 3
(cuda-gdb) n
10 X[0] = mydat[0];
(cuda-gdb) n
11 X[1] = mydat[1];
(cuda-gdb) n
12 X[2] = mydat[2];
(cuda-gdb) n
13 Y[0] = mydat[3];
(cuda-gdb) n
14 Y[1] = mydat[4];
(cuda-gdb) n
15 Y[2] = mydat[5];
(cuda-gdb) p X
$4 = (@generic double * @register) 0x1000000
(cuda-gdb) p X[0]
$5 = 0
(cuda-gdb) p X[1]
$6 = 1
(cuda-gdb) p Y[0]
$7 = 3
(cuda-gdb) p Y[1]
$8 = 4
(cuda-gdb) n
18 __syncthreads();
(cuda-gdb) n
22 for (int i = 0; i < 3; i++) {
(cuda-gdb) n
23 result += X[i] + Y[i];
(cuda-gdb) p i
$9 = 0
(cuda-gdb) p X[0]
$10 = 0
(cuda-gdb) p X[i]
Error: Failed to read global memory at address 0x0 on device 0 sm 0 warp 0 lane 0 (error=7).
??これは： 'extern __shared__ shr [];'は有効ではありませんCUDA C/C++。 –
申し訳ありません@RobertCrovella私はこれをちょっと急いで昨夜入力しました。問題を引き起こしているコードでは、 'extern __shared__ double shr [];と宣言されています。これを反映するようにOPを更新しています。 – FaceInvader
一般に、表示されていないコードの実行時の動作を説明するのは非常に難しいでしょう。 XとYのインデックス計算は間違っています（スレッドID = 0のYはスレッドID = 1のXと同じになります）が、他に何が起こっていないかは[MCVE] 。もしあなたがそれを提供できないなら、私はあなたが答えを得ることは疑うでしょう – talonmies