Introduction

備忘録。

GPU 側のコードの動きを調べるために printf を追加したけど何も表示されない事象に遭遇。
加えて、結果が常に 0 になるという不可思議な事象。

How to resove?

ブロック当たりのスレッド数が大きすぎる

まずは Stackoverflow で解決策が見つかった。

printf() in my CUDA kernel doesn’t result produce any output

下記のコードで

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
#include <iostream>

__global__ void hello_cuda()
{
    printf("Hello from CUDA kernel!\n");
}

#define BLOCK_SIZE 1
#define THREAD_PER_BLOCK 1024

int main()
{
    printf("Start!\n");

    hello_cuda<<<BLOCK_SIZE, THREAD_PER_BLOCK>>>();
    cudaDeviceSynchronize();

    printf("End!\n");
    
    return 0;
}

#define THREAD_PER_BLOCK 1024 を #define THREAD_PER_BLOCK 2048 に変えると、 hello_cuda の printf が何も表示してくれなくなる。
ややこしいのが、ホスト側の printf("End!\n"); は表示され、プログラムとしてクラッシュしないことである。

原因は GPU の性能を超えたスレッドを起動したことによるもの。
デバイスの最大ブロック数は cudaDeviceProp.maxThreadsPerBlock で確認できる。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
#include <iostream>
#include <cuda_runtime.h>

int main()
{
    int deviceCount = 0;
    cudaGetDeviceCount(&deviceCount);

    for (int i = 0; i < deviceCount; ++i)
    {
        cudaDeviceProp prop;
        cudaGetDeviceProperties(&prop, i);

        std::cout << "Device " << i << ": " << prop.name << std::endl;
        std::cout << "  Compute capability: " << prop.major << "." << prop.minor << std::endl;
        std::cout << "  Total global memory: " << prop.totalGlobalMem / (1024 * 1024) << " MB" << std::endl;
        std::cout << "  Multiprocessors: " << prop.multiProcessorCount << std::endl;
        std::cout << "  Clock rate: " << prop.clockRate / 1000 << " MHz" << std::endl;
        std::cout << "  Max threads per block: " << prop.maxThreadsPerBlock << std::endl;
        std::cout << "  Max threads dim: [" 
                  << prop.maxThreadsDim[0] << ", "
                  << prop.maxThreadsDim[1] << ", "
                  << prop.maxThreadsDim[2] << "]" << std::endl;
        std::cout << "  Max grid size: [" 
                  << prop.maxGridSize[0] << ", "
                  << prop.maxGridSize[1] << ", "
                  << prop.maxGridSize[2] << "]" << std::endl;
        std::cout << std::endl;
    }

    return 0;
}

printf が何も表示されないどころか、計算すら行われないので結果が常に 0 になっていたわけである。

CUDA_ARCHITECTURES が不一致

これは自宅で検証しているときに気付いた。
GPU の Compute Capability は下記のサイトで見つけられる。

Your GPU Compute Capability

ここで CMakeLists.txt 等で CUDA_ARCHITECTURES に、自分の GPU の Compute Capability より大きな値を設定すると、printf が表示されない。
というよりも、計算すら行われないのだろう。

例えば、自分の GPU の Compute Capability が 6.1 の時ならば CUDA_ARCHITECTURES の上限は 61 になる。
ここで、 75 を指定すると、この事象が起きてしまう。