NVIDIA OpenCL Bug Report: include directive causes CL_OUT_OF_HOST_MEMORY
Summary
Using #include directives in OpenCL kernel source code causes clBuildProgram() to fail with CL_OUT_OF_HOST_MEMORY (-6) on recent NVIDIA drivers (595.97). The build log is empty, providing no diagnostic information. The same kernel compiles successfully when the include content is inlined manually.
Environment
- GPU: NVIDIA GeForce RTX 2080 Ti (Compute Capability 7.5)
- Driver Version: 595.97
- Operating System: Microsoft Windows 11 Pro (Build 26200)
- OpenCL Version: As reported by driver
Steps to Reproduce
- Create a simple header file
test_include.h:
#define TEST_VALUE 42
- Create a kernel file
kernel_with_include.clthat uses the include:
#include "test_include.h"
__kernel void test_kernel(__global float* output) {
int id = get_global_id(0);
output[id] = (float)id + TEST_VALUE;
}
- Create an equivalent kernel
kernel_without_include.clwith the content inlined:
#define TEST_VALUE 42
__kernel void test_kernel(__global float* output) {
int id = get_global_id(0);
output[id] = (float)id + TEST_VALUE;
}
- Compile both kernels using clBuildProgram() with
-I.in build options
Expected Behavior
Both kernels should compile successfully since they are semantically equivalent.
Actual Behavior
- kernel_without_include.cl: Compiles successfully (CL_SUCCESS)
- kernel_with_include.cl: Fails with CL_OUT_OF_HOST_MEMORY (-6)
- Build log is empty
- No diagnostic information is provided
Test Program Output
NVIDIA OpenCL #include Bug Reproduction
========================================
Platform 0: NVIDIA CUDA
-> Selected
Platform 1: Intel(R) OpenCL Graphics
Device: NVIDIA GeForce RTX 2080 Ti
Driver version: 595.97
=== Testing: Kernel WITHOUT #include (should work) ===
Kernel file: kernel_without_include.cl
Kernel size: 202 bytes
Kernel source:
---
// Minimal kernel that WORKS without #include
#define TEST_VALUE 42
__kernel void test_kernel(__global float* output) {
int id = get_global_id(0);
output[id] = (float)id + TEST_VALUE;
}
---
Build options: -I.
Building...
Build result: CL_SUCCESS (0)
Build log:
(): Warning: Function test_kernel is a kernel, so overriding noinline attribute. The function may be inlined when called.
=== Testing: Kernel WITH #include (FAILS with CL_OUT_OF_HOST_MEMORY) ===
Kernel file: kernel_with_include.cl
Kernel size: 203 bytes
Kernel source:
---
// Minimal kernel that FAILS with #include
#include "test_include.h"
__kernel void test_kernel(__global float* output) {
int id = get_global_id(0);
output[id] = (float)id + TEST_VALUE;
}
---
Build options: -I.
Building...
Build result: CL_OUT_OF_HOST_MEMORY (-6)
Build log: (empty)
========================================
SUMMARY:
Kernel without #include: PASS
Kernel with #include: FAIL
========================================
BUG CONFIRMED: #include directive causes CL_OUT_OF_HOST_MEMORY
Impact
This bug prevents runtime compilation of OpenCL kernels that use #include directives. This is a common pattern for:
- Sharing type definitions between host and device code
- Modular kernel development
- Configuration headers
The only workaround is to preprocess kernels to inline all includes before passing to clBuildProgram().
Additional Notes
- The error code
CL_OUT_OF_HOST_MEMORYis misleading - this is not a memory issue but a compilation failure - The empty build log makes debugging impossible
- This appears to be a regression - the same code worked with older NVIDIA drivers
- Precompiled kernels (PTX) work correctly; only runtime source compilation is affected
- The offline NVIDIA OpenCL compiler (nvclcc) compiles the same source successfully
Workaround
Preprocess the kernel source to inline all #include directives before calling clBuildProgram().
Attachments
The following files are provided for reproduction:
test_nvidia_include_bug.cpp- Minimal C++ test programkernel_with_include.cl- Kernel that failskernel_without_include.cl- Equivalent kernel that workstest_include.h- Simple header file
Compilation Instructions
Using Visual Studio x64 Developer Command Prompt:
cl /EHsc test_nvidia_include_bug.cpp /I<path_to_opencl_headers> /link OpenCL.lib
Then run test_nvidia_include_bug.exe from the directory containing the kernel files.
nvidia_bug_report.zip (92.7 KB)