[llvm] [openmp] [OpenMP][Runtime] Handling crash with `OMP_TARGET_OFFLOAD=DISABLED` and invoking `omp_get_default_device()` (PR #171789)
Amit Tiwari via llvm-commits
llvm-commits at lists.llvm.org
Thu Dec 11 03:02:56 PST 2025
https://github.com/amitamd7 updated https://github.com/llvm/llvm-project/pull/171789
>From ead6e7e8221bc82d5e470a2a652ade91f92008f5 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Thu, 11 Dec 2025 04:39:45 -0500
Subject: [PATCH] offloading_disabled_resolve
---
.../test/api/omp_get_default_device_test.c | 37 +++++++
.../target_offload_disabled_default_device.c | 45 ++++++++
openmp/runtime/src/kmp_ftn_entry.h | 19 ++--
...p_target_offload_default_device_combined.c | 63 +++++++++++
...omp_target_offload_default_device_nested.c | 103 ++++++++++++++++++
...target_offload_default_device_operations.c | 89 +++++++++++++++
...omp_target_offload_default_device_simple.c | 47 ++++++++
...omp_target_offload_default_device_target.c | 69 ++++++++++++
...mp_target_offload_default_device_threads.c | 89 +++++++++++++++
9 files changed, 553 insertions(+), 8 deletions(-)
create mode 100644 offload/test/api/omp_get_default_device_test.c
create mode 100644 offload/test/offloading/target_offload_disabled_default_device.c
create mode 100644 openmp/runtime/test/env/omp_target_offload_default_device_combined.c
create mode 100644 openmp/runtime/test/env/omp_target_offload_default_device_nested.c
create mode 100644 openmp/runtime/test/env/omp_target_offload_default_device_operations.c
create mode 100644 openmp/runtime/test/env/omp_target_offload_default_device_simple.c
create mode 100644 openmp/runtime/test/env/omp_target_offload_default_device_target.c
create mode 100644 openmp/runtime/test/env/omp_target_offload_default_device_threads.c
diff --git a/offload/test/api/omp_get_default_device_test.c b/offload/test/api/omp_get_default_device_test.c
new file mode 100644
index 0000000000000..b1ca66cd57188
--- /dev/null
+++ b/offload/test/api/omp_get_default_device_test.c
@@ -0,0 +1,37 @@
+// RUN: %libomptarget-compile-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 |
+// %fcheck-generic
+//
+// Test omp_get_default_device() API behavior when offload is disabled
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ // Test 1: Default behavior
+ int dev1 = omp_get_default_device();
+ // CHECK: Test 1: {{0}}
+ printf("Test 1: %d\n", dev1);
+
+ // Test 2: After setting device
+ omp_set_default_device(3);
+ int dev2 = omp_get_default_device();
+ // CHECK: Test 2: {{0}}
+ printf("Test 2: %d\n", dev2);
+
+ // Test 3: Multiple sets
+ for (int i = 0; i < 5; i++) {
+ omp_set_default_device(i + 10);
+ int dev = omp_get_default_device();
+ // CHECK: Test 3.{{[0-4]}}: {{0}}
+ printf("Test 3.%d: %d\n", i, dev);
+ }
+
+ // Test 4: Consistency with initial device
+ int initial = omp_get_initial_device();
+ int default_dev = omp_get_default_device();
+ // CHECK: Test 4: EQUAL
+ printf("Test 4: %s\n", (initial == default_dev) ? "EQUAL" : "NOT_EQUAL");
+
+ return 0;
+}
diff --git a/offload/test/offloading/target_offload_disabled_default_device.c b/offload/test/offloading/target_offload_disabled_default_device.c
new file mode 100644
index 0000000000000..20b40a43bb700
--- /dev/null
+++ b/offload/test/offloading/target_offload_disabled_default_device.c
@@ -0,0 +1,45 @@
+// RUN: %libomptarget-compile-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 |
+// %fcheck-generic
+//
+// Test that setting default device before disabling offload doesn't crash
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ // Set high default device number
+ omp_set_default_device(5);
+
+ // This simulates OMP_TARGET_OFFLOAD=disabled being set after device is chosen
+ // In practice, the environment variable is read at runtime init
+
+ // CHECK: num_devices: 0
+ printf("num_devices: %d\n", omp_get_num_devices());
+
+ // CHECK: initial_device: 0
+ printf("initial_device: %d\n", omp_get_initial_device());
+
+ // CHECK: default_device: 0
+ printf("default_device: %d\n", omp_get_default_device());
+
+ // Target region should execute on host
+ int result = -1;
+#pragma omp target map(from : result)
+ {
+ result = omp_get_device_num();
+ }
+
+ // CHECK: executed_on: 0
+ printf("executed_on: %d\n", result);
+
+ // CHECK: PASS
+ if (result == omp_get_initial_device() &&
+ omp_get_default_device() == omp_get_initial_device()) {
+ printf("PASS\n");
+ return 0;
+ }
+
+ printf("FAIL\n");
+ return 1;
+}
diff --git a/openmp/runtime/src/kmp_ftn_entry.h b/openmp/runtime/src/kmp_ftn_entry.h
index 6adf60e7ad210..4e7c2822b9358 100644
--- a/openmp/runtime/src/kmp_ftn_entry.h
+++ b/openmp/runtime/src/kmp_ftn_entry.h
@@ -1143,14 +1143,6 @@ int FTN_STDCALL KMP_EXPAND_NAME(FTN_GET_TEAM_NUM)(void) {
#endif
}
-int FTN_STDCALL KMP_EXPAND_NAME(FTN_GET_DEFAULT_DEVICE)(void) {
-#if KMP_MIC || KMP_OS_DARWIN || defined(KMP_STUB)
- return 0;
-#else
- return __kmp_entry_thread()->th.th_current_task->td_icvs.default_device;
-#endif
-}
-
void FTN_STDCALL KMP_EXPAND_NAME(FTN_SET_DEFAULT_DEVICE)(int KMP_DEREF arg) {
#if KMP_MIC || KMP_OS_DARWIN || defined(KMP_STUB)
// Nothing.
@@ -1197,6 +1189,17 @@ int FTN_STDCALL KMP_EXPAND_NAME(FTN_GET_INITIAL_DEVICE)(void) {
return KMP_EXPAND_NAME(FTN_GET_NUM_DEVICES)();
}
+int FTN_STDCALL KMP_EXPAND_NAME(FTN_GET_DEFAULT_DEVICE)(void) {
+#if KMP_MIC || KMP_OS_DARWIN || defined(KMP_STUB)
+ return 0;
+#else
+ // When offloading is disabled, return the initial device (host)
+ if (__kmp_target_offload == tgt_disabled)
+ return KMP_EXPAND_NAME(FTN_GET_INITIAL_DEVICE)();
+ return __kmp_entry_thread()->th.th_current_task->td_icvs.default_device;
+#endif
+}
+
#if defined(KMP_STUB)
// Entries for stubs library
// As all *target* functions are C-only parameters always passed by value
diff --git a/openmp/runtime/test/env/omp_target_offload_default_device_combined.c b/openmp/runtime/test/env/omp_target_offload_default_device_combined.c
new file mode 100644
index 0000000000000..4bd8c57bc5785
--- /dev/null
+++ b/openmp/runtime/test/env/omp_target_offload_default_device_combined.c
@@ -0,0 +1,63 @@
+// RUN: %libomp-compile-and-run
+// REQUIRES: ompt
+//
+// Test that omp_get_default_device() returns the initial device (0) when
+// OMP_TARGET_OFFLOAD=DISABLED, with both OMP_DEFAULT_DEVICE environment
+// variable and omp_set_default_device() API call setting non-zero values.
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <omp.h>
+
+extern void kmp_set_defaults(char const *str);
+
+int main() {
+ // Simulate worst case: both environment variable and API call set non-zero
+ // device
+ kmp_set_defaults("OMP_DEFAULT_DEVICE=3");
+ kmp_set_defaults("OMP_TARGET_OFFLOAD=DISABLED");
+
+// Force parallel region to initialize runtime
+#pragma omp parallel
+ {
+ }
+
+ int initial_device = omp_get_initial_device();
+ int default_device_1 = omp_get_default_device();
+
+ printf("With OMP_DEFAULT_DEVICE=3 and OMP_TARGET_OFFLOAD=DISABLED:\n");
+ printf(" initial_device = %d\n", initial_device);
+ printf(" default_device = %d\n", default_device_1);
+
+ if (default_device_1 != initial_device) {
+ fprintf(stderr,
+ "FAIL: Environment variable not overridden by offload disabled\n");
+ return EXIT_FAILURE;
+ }
+
+ // Now also call omp_set_default_device()
+ omp_set_default_device(7);
+ int default_device_2 = omp_get_default_device();
+
+ printf("After additional omp_set_default_device(7):\n");
+ printf(" default_device = %d\n", default_device_2);
+
+ if (default_device_2 != initial_device) {
+ fprintf(stderr, "FAIL: API call not overridden by offload disabled\n");
+ return EXIT_FAILURE;
+ }
+
+ // Verify consistency across multiple calls
+ for (int i = 0; i < 5; i++) {
+ int dev = omp_get_default_device();
+ if (dev != initial_device) {
+ fprintf(stderr,
+ "FAIL: Inconsistent result on call %d: got %d, expected %d\n", i,
+ dev, initial_device);
+ return EXIT_FAILURE;
+ }
+ }
+
+ printf("PASS: default_device consistently returns initial_device\n");
+ return EXIT_SUCCESS;
+}
diff --git a/openmp/runtime/test/env/omp_target_offload_default_device_nested.c b/openmp/runtime/test/env/omp_target_offload_default_device_nested.c
new file mode 100644
index 0000000000000..6ad120e0bff80
--- /dev/null
+++ b/openmp/runtime/test/env/omp_target_offload_default_device_nested.c
@@ -0,0 +1,103 @@
+// RUN: %libomp-compile-and-run
+// REQUIRES: ompt
+//
+// Test that omp_get_default_device() returns the initial device consistently
+// across nested parallel regions and with ICV inheritance when
+// OMP_TARGET_OFFLOAD=DISABLED.
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <omp.h>
+
+extern void kmp_set_defaults(char const *str);
+
+int check_default_device(const char *context, int expected_initial) {
+ int default_dev = omp_get_default_device();
+ int initial_dev = omp_get_initial_device();
+
+ if (default_dev != initial_dev) {
+ fprintf(stderr, "FAIL [%s]: default=%d, initial=%d\n", context, default_dev,
+ initial_dev);
+ return 1;
+ }
+
+ if (initial_dev != expected_initial) {
+ fprintf(stderr, "FAIL [%s]: initial=%d, expected=%d\n", context,
+ initial_dev, expected_initial);
+ return 1;
+ }
+
+ return 0;
+}
+
+int main() {
+ int errors = 0;
+
+ // Set configuration
+ kmp_set_defaults("OMP_DEFAULT_DEVICE=8");
+ kmp_set_defaults("OMP_TARGET_OFFLOAD=DISABLED");
+
+// Initialize runtime
+#pragma omp parallel
+ {
+ }
+
+ int initial_device = omp_get_initial_device();
+ printf("initial_device = %d\n", initial_device);
+
+ // Test 1: Sequential region
+ errors += check_default_device("sequential", initial_device);
+
+// Test 2: Parallel region
+#pragma omp parallel reduction(+ : errors)
+ {
+ errors += check_default_device("parallel", initial_device);
+
+// Test 3: Nested parallel (if supported)
+#pragma omp parallel reduction(+ : errors) if (omp_get_max_threads() > 2)
+ {
+ errors += check_default_device("nested parallel", initial_device);
+ }
+ }
+
+// Test 4: After modifying in one thread
+#pragma omp parallel num_threads(4) reduction(+ : errors)
+ {
+ int tid = omp_get_thread_num();
+
+ // Each thread tries to set different default device
+ omp_set_default_device(tid + 20);
+
+ // But should still get initial device
+ errors += check_default_device("after thread-local set", initial_device);
+
+#pragma omp barrier
+
+ // Check again after barrier
+ errors += check_default_device("after barrier", initial_device);
+ }
+
+ // Test 5: Back in sequential after all the parallel regions
+ errors += check_default_device("sequential final", initial_device);
+
+ // Test 6: Target region context
+ int target_errors = 0;
+#pragma omp target map(tofrom : target_errors)
+ {
+ int default_dev = omp_get_default_device();
+ int initial_dev = omp_get_initial_device();
+ if (default_dev != initial_dev) {
+ target_errors = 1;
+ }
+ }
+ errors += target_errors;
+
+ if (errors > 0) {
+ fprintf(stderr, "FAIL: %d errors detected\n", errors);
+ return EXIT_FAILURE;
+ }
+
+ printf("PASS: default_device consistently returns initial_device across all "
+ "contexts\n");
+ return EXIT_SUCCESS;
+}
diff --git a/openmp/runtime/test/env/omp_target_offload_default_device_operations.c b/openmp/runtime/test/env/omp_target_offload_default_device_operations.c
new file mode 100644
index 0000000000000..6799b74a67349
--- /dev/null
+++ b/openmp/runtime/test/env/omp_target_offload_default_device_operations.c
@@ -0,0 +1,89 @@
+// RUN: %libomp-compile-and-run
+// REQUIRES: ompt
+//
+// Test that device operations using omp_get_default_device() don't crash
+// when OMP_TARGET_OFFLOAD=DISABLED. This simulates real-world usage where
+// the default device is used for device-specific operations.
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <omp.h>
+
+extern void kmp_set_defaults(char const *str);
+
+int main() {
+ // Simulate the problematic scenario: high default device number + disabled
+ // offload Use API call instead of env var to ensure ICV is set
+ omp_set_default_device(10);
+
+ // Now disable offload
+ kmp_set_defaults("OMP_TARGET_OFFLOAD=DISABLED");
+
+// Force parallel region to initialize runtime
+#pragma omp parallel
+ {
+ }
+
+ int device = omp_get_default_device();
+ int num_devices = omp_get_num_devices();
+ int initial_device = omp_get_initial_device();
+
+ printf("Configuration:\n");
+ printf(" num_devices = %d\n", num_devices);
+ printf(" initial_device = %d\n", initial_device);
+ printf(" default_device = %d\n", device);
+
+ // Verify device is in valid range
+ if (device < 0 || device > num_devices) {
+ fprintf(stderr, "FAIL: default_device (%d) is out of valid range [0, %d]\n",
+ device, num_devices);
+ return EXIT_FAILURE;
+ }
+
+ // Test 1: Check if we're on initial device (should be true when offload
+ // disabled)
+ int is_initial = (device == initial_device);
+ if (!is_initial) {
+ fprintf(stderr,
+ "FAIL: default_device (%d) is not the initial_device (%d)\n",
+ device, initial_device);
+ return EXIT_FAILURE;
+ }
+
+ // Test 2: Use device in target region with device clause
+ // This should not crash even though OMP_DEFAULT_DEVICE=10
+ int result = -1;
+#pragma omp target device(device) map(from : result)
+ {
+ result = omp_get_device_num();
+ }
+
+ printf("Target region executed on device: %d\n", result);
+
+ // When offload is disabled, target should execute on host (initial device)
+ if (result != initial_device) {
+ fprintf(stderr, "FAIL: Target executed on device %d, expected %d\n", result,
+ initial_device);
+ return EXIT_FAILURE;
+ }
+
+ // Test 3: Query device properties using the default device
+ int is_host = (device == initial_device);
+ printf("Device %d is_host: %d\n", device, is_host);
+
+ // Test 4: Verify target region with default device specification
+ int test_value = 0;
+#pragma omp target device(omp_get_default_device()) map(tofrom : test_value)
+ {
+ test_value = 42;
+ }
+
+ if (test_value != 42) {
+ fprintf(stderr,
+ "FAIL: Target region with default device did not execute\n");
+ return EXIT_FAILURE;
+ }
+
+ printf("PASS: All device operations completed without crash\n");
+ return EXIT_SUCCESS;
+}
diff --git a/openmp/runtime/test/env/omp_target_offload_default_device_simple.c b/openmp/runtime/test/env/omp_target_offload_default_device_simple.c
new file mode 100644
index 0000000000000..a4a5479e154b3
--- /dev/null
+++ b/openmp/runtime/test/env/omp_target_offload_default_device_simple.c
@@ -0,0 +1,47 @@
+// RUN: %libomp-compile-and-run
+// REQUIRES: ompt
+//
+// Simple smoke test to verify omp_get_default_device() returns initial device
+// when OMP_TARGET_OFFLOAD=DISABLED with OMP_DEFAULT_DEVICE=2.
+// This is the C equivalent of the Fortran smoke test.
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <omp.h>
+
+extern void kmp_set_defaults(char const *str);
+
+int main() {
+ // Key to reproducing bug: Set default device BEFORE disabling offload
+ // This ensures the ICV contains a non-zero value
+ omp_set_default_device(2);
+
+ // Now disable offload
+ kmp_set_defaults("OMP_TARGET_OFFLOAD=DISABLED");
+
+// Initialize runtime
+#pragma omp parallel
+ {
+ }
+
+ int num_devices = omp_get_num_devices();
+ int initial_device = omp_get_initial_device();
+ int default_device = omp_get_default_device();
+
+ // Print results
+ printf("number of devices %d\n", num_devices);
+ printf("initial device %d\n", initial_device);
+ printf("default device %d\n", default_device);
+
+ // The key test: default device should equal initial device
+ if (initial_device == default_device) {
+ printf("PASS\n");
+ return EXIT_SUCCESS;
+ } else {
+ fprintf(stderr, "FAIL: default_device (%d) != initial_device (%d)\n",
+ default_device, initial_device);
+ fprintf(stderr, "This would cause: device number '%d' out of range\n",
+ default_device);
+ return EXIT_FAILURE;
+ }
+}
diff --git a/openmp/runtime/test/env/omp_target_offload_default_device_target.c b/openmp/runtime/test/env/omp_target_offload_default_device_target.c
new file mode 100644
index 0000000000000..a821a0547dda5
--- /dev/null
+++ b/openmp/runtime/test/env/omp_target_offload_default_device_target.c
@@ -0,0 +1,69 @@
+// RUN: %libomp-compile-and-run
+// REQUIRES: ompt
+//
+// Test that omp_get_default_device() returns the initial device (0) when
+// called from within a target region when OMP_TARGET_OFFLOAD=DISABLED.
+// The target region should execute on the host.
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <omp.h>
+
+extern void kmp_set_defaults(char const *str);
+
+int main() {
+ // Set non-zero default device using API (more direct than env var)
+ omp_set_default_device(4);
+
+ // Now disable offload
+ kmp_set_defaults("OMP_TARGET_OFFLOAD=DISABLED");
+
+// Force parallel region to initialize runtime
+#pragma omp parallel
+ {
+ }
+
+ int initial_device = omp_get_initial_device();
+ int host_default_device = omp_get_default_device();
+ int target_default_device = -1;
+ int target_is_initial = -1;
+
+ printf("Host context:\n");
+ printf(" initial_device = %d\n", initial_device);
+ printf(" default_device = %d\n", host_default_device);
+
+// Call omp_get_default_device() from within target region
+// When offload is disabled, this should execute on host
+#pragma omp target map(from : target_default_device, target_is_initial)
+ {
+ target_default_device = omp_get_default_device();
+ target_is_initial = omp_is_initial_device();
+ }
+
+ printf("Target context (executed on host when offload disabled):\n");
+ printf(" default_device = %d\n", target_default_device);
+ printf(" is_initial_device = %d\n", target_is_initial);
+
+ // When offload is disabled, target region executes on host
+ if (target_is_initial != 1) {
+ fprintf(stderr, "FAIL: Target region did not execute on initial device\n");
+ return EXIT_FAILURE;
+ }
+
+ // Both host and target context should return same device
+ if (host_default_device != initial_device) {
+ fprintf(stderr, "FAIL: Host default_device (%d) != initial_device (%d)\n",
+ host_default_device, initial_device);
+ return EXIT_FAILURE;
+ }
+
+ if (target_default_device != initial_device) {
+ fprintf(stderr, "FAIL: Target default_device (%d) != initial_device (%d)\n",
+ target_default_device, initial_device);
+ return EXIT_FAILURE;
+ }
+
+ printf("PASS: default_device returns initial_device in both host and target "
+ "contexts\n");
+ return EXIT_SUCCESS;
+}
diff --git a/openmp/runtime/test/env/omp_target_offload_default_device_threads.c b/openmp/runtime/test/env/omp_target_offload_default_device_threads.c
new file mode 100644
index 0000000000000..1640f6366f86c
--- /dev/null
+++ b/openmp/runtime/test/env/omp_target_offload_default_device_threads.c
@@ -0,0 +1,89 @@
+// RUN: %libomp-compile-and-run
+// REQUIRES: ompt
+//
+// Test that omp_get_default_device() consistently returns the initial device
+// across multiple threads when OMP_TARGET_OFFLOAD=DISABLED.
+// This ensures thread-safety.
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <omp.h>
+
+extern void kmp_set_defaults(char const *str);
+
+int main() {
+ const int NUM_THREADS = 8;
+ const int NUM_ITERATIONS = 100;
+
+ // Set non-zero default device and disable offload
+ kmp_set_defaults("OMP_DEFAULT_DEVICE=6");
+ kmp_set_defaults("OMP_TARGET_OFFLOAD=DISABLED");
+
+// Force parallel region to initialize runtime
+#pragma omp parallel
+ {
+ }
+
+ int initial_device = omp_get_initial_device();
+ int errors = 0;
+
+ printf("Testing with %d threads, %d iterations each\n", NUM_THREADS,
+ NUM_ITERATIONS);
+ printf("initial_device = %d\n", initial_device);
+
+// Test across multiple parallel regions and threads
+#pragma omp parallel num_threads(NUM_THREADS) reduction(+ : errors)
+ {
+ int tid = omp_get_thread_num();
+
+ for (int i = 0; i < NUM_ITERATIONS; i++) {
+ int default_device = omp_get_default_device();
+
+ if (default_device != initial_device) {
+#pragma omp critical
+ {
+ fprintf(
+ stderr,
+ "FAIL: Thread %d iteration %d: default_device=%d, expected=%d\n",
+ tid, i, default_device, initial_device);
+ }
+ errors++;
+ }
+
+ // Also test after setting default device in each thread
+ if (i % 10 == 0) {
+ omp_set_default_device(tid + 10);
+ default_device = omp_get_default_device();
+
+ if (default_device != initial_device) {
+#pragma omp critical
+ {
+ fprintf(
+ stderr,
+ "FAIL: Thread %d after set: default_device=%d, expected=%d\n",
+ tid, default_device, initial_device);
+ }
+ errors++;
+ }
+ }
+ }
+ }
+
+ if (errors > 0) {
+ fprintf(stderr, "FAIL: %d errors detected across all threads\n", errors);
+ return EXIT_FAILURE;
+ }
+
+ // Final verification
+ int final_device = omp_get_default_device();
+ if (final_device != initial_device) {
+ fprintf(stderr,
+ "FAIL: Final check failed: default_device=%d, expected=%d\n",
+ final_device, initial_device);
+ return EXIT_FAILURE;
+ }
+
+ printf("PASS: default_device consistently returns initial_device across all "
+ "threads\n");
+ return EXIT_SUCCESS;
+}
More information about the llvm-commits
mailing list