helloTile: per-sample build fails on Ubuntu/Debian/Fedora — _FORTIFY_SOURCE rewrites printf to a non-tile-callable __printf_chk (v13.3)

[yashashav-dk]

Summary

Building the cpp/9_CUDA_Tile/helloTile sample on its own (the per-sample build the README documents — "you can also follow this process … from within any individual sample") fails from a clean v13.3 checkout on a default Ubuntu 24.04 system, and on any distro whose GCC package defaults -D_FORTIFY_SOURCE on at -O1+ (Debian, Fedora). helloTile.cu calls printf from a __tile_global__ kernel, which CUDA 13.3 intends to allow — but with FORTIFY active, printf is rewritten to __printf_chk, which is not tile-callable, so nvcc rejects the call.

The whole-tree build from the repo root works (the FORTIFY workaround in the parent 9_CUDA_Tile/CMakeLists.txt is inherited there), so this is specific to the documented per-sample build path.

Environment

Component	Value
OS	Ubuntu 24.04.4 LTS (also affects Debian / Fedora; vanilla-GCC distros like Arch are unaffected)
CUDA Toolkit	13.3.r13.3 (nvcc V13.3.33)
GCC	13.3.0 (Ubuntu 13.3.0-6ubuntu2~24.04.1)
glibc	2.39
CMake	4.3.2 (also reproduces with 3.28.3)
cuda-samples	v13.3 (b7c5481)
GPU	RTX PRO 6000 Blackwell (sm_120)

Reproduction

Per-sample build (documented as supported), which fails:

git clone --depth 1 --branch v13.3 https://github.com/NVIDIA/cuda-samples
cd cuda-samples/cpp/9_CUDA_Tile/helloTile
mkdir build && cd build
cmake .. -DCMAKE_CUDA_ARCHITECTURES="120" -DCMAKE_BUILD_TYPE=Release
make

Actual

helloTile.cu(48): error: calling a __host__ __device__ function("printf")
                   from a __tile_global__ function("tileKernel") is not allowed
helloTile.cu(49): error: calling a __host__ __device__ function("printf")
                   from a __tile_global__ function("tileKernel") is not allowed
helloTile.cu(52): error: calling a __host__ __device__ function("printf")
                   from a __tile_global__ function("tileKernel") is not allowed
3 errors detected in the compilation of "helloTile.cu".

Expected

The per-sample build succeeds (as the whole-tree build does) and the binary prints the Hello, SIMT! / Hello, Tile! / Hello, Host! sequence.

For contrast, the whole-tree build from the repo root does build this sample correctly:

cd cuda-samples
cmake -S . -B build -DCMAKE_CUDA_ARCHITECTURES="120" -DCMAKE_BUILD_TYPE=Release
cmake --build build --target helloTile   # succeeds

Root cause

1. CUDA 13.3 declares printf as tile-callable. crt/common_functions.h re-declares printf with the tile / tile_builtin attributes (via the __NV_TL__ / __NV_TL_BUILTIN__ macros), so calling it from a __tile_global__ kernel is intended to be legal.
2. GCC enables FORTIFY by default at -O1+. On Ubuntu/Debian/Fedora the GCC package injects -D_FORTIFY_SOURCE for optimized builds. On this box:

   $ echo | gcc -O2 -dM -E - | grep FORTIFY
   #define _FORTIFY_SOURCE 3
   # still present with -nostdinc, confirming GCC injects it (not a glibc header)
   

   The samples Release build uses -O2, so FORTIFY is active.
3. glibc then rewrites printf → __printf_chk. Under __USE_FORTIFY_LEVEL > 0, <bits/stdio2.h> provides an inline overload and a #define printf(...) macro that forward to __printf_chk, which glibc declares with no CUDA/tile attributes. A __tile_global__ kernel may only call __device__/__tile__ functions, so the redirected call is rejected.

(The diagnostic mentions printf because that's the original declaration; the actual call target after FORTIFY rewriting is the non-tile __printf_chk.)

Why only the per-sample build is affected

cpp/9_CUDA_Tile/CMakeLists.txt (the parent) already contains the fix:

add_compile_options(
    $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=-U_FORTIFY_SOURCE>
    $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=-D_FORTIFY_SOURCE=0>
)

add_compile_options sets a directory property. It is inherited when helloTile is built through that parent (the whole-tree build from the repo root), but not when helloTile/ is configured as its own top-level project — which is exactly the documented per-sample path. So the single-sample build a user would naturally run is the one that misses the workaround.

Note: building from the category directory (cmake -S cpp/9_CUDA_Tile) is not an alternative either — that parent has no project() / cmake_minimum_required(), so it can't be configured as a standalone top-level build.

Suggested fix

Minimal (per-leaf): add the same workaround to cpp/9_CUDA_Tile/helloTile/CMakeLists.txt, after find_package(CUDAToolkit REQUIRED):

# GCC packages (Ubuntu/Debian/Fedora) default _FORTIFY_SOURCE on at -O1+, which
# rewrites printf -> __printf_chk; __printf_chk is not callable from tile kernels.
# Needed per-leaf because the parent dir's add_compile_options is not inherited
# when this sample is configured as its own top-level project (per-sample build).
add_compile_options(
    $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=-U_FORTIFY_SOURCE>
    $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=-D_FORTIFY_SOURCE=0>
)

A cleaner variant: move the workaround into a shared cmake/Modules/ helper and include() it from each 9_CUDA_Tile/* leaf, so it survives leaf-only configures and covers any future tile sample that prints.

Verified: with this change the per-sample build succeeds and runs correctly on Blackwell (sm_120):

Hello, SIMT!
[SIMT] *x == 0
[SIMT] *x = 100

Hello, Tile!
[Tile] *x == 100
[Tile] *x = 200

Hello, Host!
[Host] *x == 200

(-O0 and -Xcompiler=-U_FORTIFY_SOURCE -Xcompiler=-D_FORTIFY_SOURCE=0 on the command line also work as ad-hoc workarounds.)

Separate, unrelated note

While checking the configure roots, a distinct v13.3 issue surfaced: cmake -S cpp (configuring the C++ subtree directly, rather than from the repo root) fails because several 4_CUDA_Libraries/* leaves declare project(<name> LANGUAGES CXX) but then call target_compile_features(... cuda_std_17) (no known features for CUDA compiler ""). It doesn't affect the repo-root build (which enables CUDA globally) and is independent of the tile/printf issue above — happy to file it separately; mentioned only for completeness.