0

何らかの理由で、特定のカーネルに設定したブレークポイントが完全に無視されます... でエラー ステータスを確認しcudaGetLastError()たところ、すべて正常に実行されたことがわかりました。これは、カーネルが実行されたことを意味していると確信しています。ステートメントを配置printfしても、何も出力されないため、余分な情報は得られません。デバッグモードに入ったカーネルでも、printf呼び出しは効果がありません。ここで何が問題になる可能性がありますか?!

Tesla M2075 (ドライバー バージョン 295.41) で Cuda 4.2 を実行しています。デバッグ時の出力:

(cuda-gdb) break cudaCalcBeamIntersect
Breakpoint 1 at 0x401cfb: file cudacalcbeamintersect.cu, line 109.
(cuda-gdb) r
Starting program: /home/heit/cuda/vfind/vfind singleevent.txt 1 1 1 
[Thread debugging using libthread_db enabled]
[New Thread 0x7ffff5dd5700 (LWP 20241)]
[Context Create of context 0x634220 on Device 0]
[Launch of CUDA Kernel 0 (memset32_post<<<(64,1,1),(64,1,1)>>>) on Device 0]
[Launch of CUDA Kernel 1 (memset32_post<<<(8,1,1),(64,1,1)>>>) on Device 0]
[Launch of CUDA Kernel 2 (memset32_post<<<(64,1,1),(64,1,1)>>>) on Device 0]
[Launch of CUDA Kernel 3 (memset32_post<<<(1,1,1),(64,1,1)>>>) on Device 0]
[Launch of CUDA Kernel 4 (memset32_post<<<(1,1,1),(64,1,1)>>>) on Device 0]
[Launch of CUDA Kernel 5 (memset32_post<<<(8,1,1),(64,1,1)>>>) on Device 0]
[Launch of CUDA Kernel 6 (cudaInitializeGlobals<<<(256,1,1),(128,1,1)>>>) on Device 0]
no error
[Launch of CUDA Kernel 7 (cudaCalcBeamIntersect<<<(256,1,1),(128,1,1)>>>) on Device 0]
no error
Elapsed time: 0.876842 seconds.
[Thread 0x7ffff5dd5700 (LWP 20241) exited]
[Termination of CUDA Kernel 6 (cudaInitializeGlobals<<<(256,1,1),(128,1,1)>>>) on Device 0]

Program exited normally.

「エラーなし」の出力は、 を呼び出すことによってカーネルの外側に出力されcout << cudaGetErrorString(cudaGetLastError()) << '\n';、cudaInitializeGlobals()(cuda-gdb でステップスルーできる) と のcudaCalcBeamIntersect()両方が問題なく実行されることを示します。ただし、後者はデバッグできません。

問題のカーネルはまだ暫定的なものであり、(静的) グローバル メモリに格納されるいくつかの値を計算します。これらの値に対して他に何も行われないので、コンパイラがこの呼び出しを完全に最適化する可能性はありますか? もしそうなら、なぜですか？そして、この動作を防ぐ方法は?? (-O0 は無効)

乾杯！

編集 - コード:

** カーネルを呼び出すコード **

    uint const nEvents = events.size();     // total number of events

    /* Not important ... */

// Allocate memory to hold the events
    Track *dev_events;                      
    cudaMalloc(&dev_events, linearEvents.size() * sizeof(Track));

// Copy all events to the GPU
    cudaMemcpy(dev_events, &linearEvents[0], linearEvents.size() * sizeof(Track), cudaMemcpyHostToDevice);

// Initialize the global data, like the histogram and the array of z-values
    cudaInitializeGlobals <<< tpb, bpg >>> ();
    cout << cudaGetErrorString(cudaGetLastError()) << '\n';

    cout << "Processing " << nEvents << " event(s)\n";
    uint linearIdx = 0;
    for (uint event = 0; event != nEvents; ++event)
    {
        uint nTracks = events[event].size();

        if (nTracks > MAX_NUMBER_OF_TRACKS)
        {
            cout << "Number of tracks in event " << event << " exceeds maximum number of tracks.\n";
            exit(1);
        }

        cudaCalcBeamIntersect <<< tpb, bpg >>> (dev_events + linearIdx, nTracks, bipThresh, binWidth);
        cout << cudaGetErrorString(cudaGetLastError()) << '\n';

    // Update linear index
        linearIdx += nTracks;
    }

cudacalcbeamintersect.cu

#include "vfind.cuh"

__device__ float    dev_zMin;
__device__ float    dev_zMax;
__device__ float    dev_zValues[MAX_NUMBER_OF_TRACKS];
__device__ uint     dev_histogram[MAX_NUMBER_OF_BINS];

__constant__ Track dev_beam = 
{
    {0, 0, 1},
    {0, 0, 0}
};

__global__ void cudaInitializeGlobals()
{
    uint const tid = threadIdx.x + blockIdx.x * blockDim.x;
    uint const nThreads = blockDim.x * gridDim.x;

    if (tid == 0)
    {
        dev_zMin = 1e6;
        dev_zMax = -1e6;
    }

    uint idx = tid;
    while (idx < MAX_NUMBER_OF_BINS || idx < MAX_NUMBER_OF_TRACKS)          
    {
        if (idx < MAX_NUMBER_OF_BINS)
            dev_histogram[idx] = 0;

        if (idx < MAX_NUMBER_OF_TRACKS)
            dev_zValues[idx] = 0;

        idx += nThreads;
    }
}

__device__ float dot(float const v1[3], float const v2[3])
{
    // Stuff
}

__device__ float distance(Track const &t1, Track const &t2)
{
    // Even more boring unimportant stuff
}

__device__ Vertex vertex(Track const &t1, Track const &t2)
{
    // Yet even more boring unimportant stuff
}

__global__ void cudaCalcBeamIntersect(Track const *tracks, uint nTracks, float bipTresh, float binWidth)
{
    uint const tid = threadIdx.x + blockIdx.x * blockDim.x;
    uint const nThreads = blockDim.x * gridDim.x;

    uint idx = tid;
    while (idx < nTracks)
    {
        float dist = distance(tracks[idx], dev_beam);
        if (dist < bipTresh)
        {
            float z = vertex(tracks[idx], dev_beam).z;

            if (z < dev_zMin)
                atomicExch(&dev_zMin, z);

            if (z > dev_zMax)
                atomicExch(&dev_zMax, z);

            dev_zValues[idx] = z;
        }

        idx += nThreads;
    }

    __syncthreads();

    // To be continued here
}
4

1 に答える 1

1

@JorenHeitカーネルcudaCalcBeamIntersectにはグローバルメモリの副作用があり、最適化されるべきではありません。投稿されたcuda-gdbの出力に基づくと、作業を開始したホストスレッドは、作業の完了を待機していないようです（cudaDeviceSynchronize()呼び出しまたはcudaMemcpyデバイスからホストへ）。その結果、cudaCalcBeamIntersect kernelGPUで実行される前に、ホストスレッドが終了します。cudaDeviceSynchronize()アプリケーションでカーネルを起動するたびに、呼び出しを追加してみてください。

于 2013-02-19T18:24:01.490 に答える