[llvm] [openmp] [OpenMP][Runtime] Handling crash with OMP_TARGET_OFFLOAD=DISABLED and invalid default device (PR #173630)
Amit Tiwari via llvm-commits
llvm-commits at lists.llvm.org
Tue Dec 30 00:53:42 PST 2025
https://github.com/amitamd7 updated https://github.com/llvm/llvm-project/pull/173630
>From b7ca88d2cbd3d4232b71fddc7d546ef107518ff2 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Fri, 26 Dec 2025 05:57:27 -0500
Subject: [PATCH] default_device_setter_resolution
---
.../api/omp_set_default_device_disabled.c | 53 +++++++++++
..._offload_disabled_set_default_device_api.c | 65 ++++++++++++++
..._offload_disabled_set_default_device_env.c | 45 ++++++++++
...fload_disabled_set_default_device_stress.c | 85 ++++++++++++++++++
openmp/runtime/src/kmp_ftn_entry.h | 42 +++++----
openmp/runtime/src/kmp_runtime.cpp | 17 +++-
...rget_offload_disabled_set_default_device.c | 72 +++++++++++++++
...oad_disabled_set_default_device_combined.c | 70 +++++++++++++++
..._offload_disabled_set_default_device_icv.c | 80 +++++++++++++++++
...fload_disabled_set_default_device_nested.c | 87 +++++++++++++++++++
...load_disabled_set_default_device_threads.c | 69 +++++++++++++++
11 files changed, 667 insertions(+), 18 deletions(-)
create mode 100644 offload/test/api/omp_set_default_device_disabled.c
create mode 100644 offload/test/offloading/target_offload_disabled_set_default_device_api.c
create mode 100644 offload/test/offloading/target_offload_disabled_set_default_device_env.c
create mode 100644 offload/test/offloading/target_offload_disabled_set_default_device_stress.c
create mode 100644 openmp/runtime/test/env/omp_target_offload_disabled_set_default_device.c
create mode 100644 openmp/runtime/test/env/omp_target_offload_disabled_set_default_device_combined.c
create mode 100644 openmp/runtime/test/env/omp_target_offload_disabled_set_default_device_icv.c
create mode 100644 openmp/runtime/test/env/omp_target_offload_disabled_set_default_device_nested.c
create mode 100644 openmp/runtime/test/env/omp_target_offload_disabled_set_default_device_threads.c
diff --git a/offload/test/api/omp_set_default_device_disabled.c b/offload/test/api/omp_set_default_device_disabled.c
new file mode 100644
index 0000000000000..e226200581a77
--- /dev/null
+++ b/offload/test/api/omp_set_default_device_disabled.c
@@ -0,0 +1,53 @@
+// RUN: %libomptarget-compile-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 |
+// %fcheck-generic
+//
+// API test: omp_set_default_device() should clamp to initial_device
+// when OMP_TARGET_OFFLOAD=disabled
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ // Force runtime initialization
+#pragma omp parallel
+ {
+ }
+
+ int initial = omp_get_initial_device();
+ int num_devices = omp_get_num_devices();
+
+ // CHECK: num_devices: 0
+ printf("num_devices: %d\n", num_devices);
+ // CHECK: initial_device: [[INITIAL:[0-9]+]]
+ printf("initial_device: %d\n", initial);
+
+ // Test 1: Set to high value
+ omp_set_default_device(10);
+ int dev1 = omp_get_default_device();
+ // CHECK: After set(10): [[INITIAL]]
+ printf("After set(10): %d\n", dev1);
+
+ // Test 2: Set to another high value
+ omp_set_default_device(5);
+ int dev2 = omp_get_default_device();
+ // CHECK: After set(5): [[INITIAL]]
+ printf("After set(5): %d\n", dev2);
+
+ // Test 3: All should be equal to initial
+ // CHECK: All equal: YES
+ printf("All equal: %s\n",
+ (dev1 == initial && dev2 == initial) ? "YES" : "NO");
+
+ // Test 4: Target region should work
+ int executed = 0;
+#pragma omp target map(tofrom : executed)
+ {
+ executed = 1;
+ }
+
+ // CHECK: Target executed: YES
+ printf("Target executed: %s\n", executed ? "YES" : "NO");
+
+ return 0;
+}
diff --git a/offload/test/offloading/target_offload_disabled_set_default_device_api.c b/offload/test/offloading/target_offload_disabled_set_default_device_api.c
new file mode 100644
index 0000000000000..e2e60293109c3
--- /dev/null
+++ b/offload/test/offloading/target_offload_disabled_set_default_device_api.c
@@ -0,0 +1,65 @@
+// RUN: %libomptarget-compile-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 |
+// %fcheck-generic
+//
+// Integration test: omp_set_default_device() should not cause crashes
+// and target regions should work when offload is disabled
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ // Force runtime initialization
+#pragma omp parallel
+ {
+ }
+
+ int initial_device = omp_get_initial_device();
+
+ // CHECK: initial_device: [[INITIAL:[0-9]+]]
+ printf("initial_device: %d\n", initial_device);
+
+ // Explicitly set high device number that would be invalid
+ omp_set_default_device(10);
+
+ int default_device = omp_get_default_device();
+
+ // Should return initial_device (not 10)
+ // CHECK: default_device: [[INITIAL]]
+ printf("default_device: %d\n", default_device);
+
+ // Verify target region works with device clause using default device
+ int executed = 0;
+#pragma omp target device(omp_get_default_device()) map(tofrom : executed)
+ {
+ executed = 1;
+ }
+
+ // CHECK: Target with device clause: PASS
+ printf("Target with device clause: %s\n", executed ? "PASS" : "FAIL");
+
+ // Try different device numbers
+ for (int i = 0; i < 5; i++) {
+ omp_set_default_device(i * 10);
+ int dev = omp_get_default_device();
+ if (dev != initial_device) {
+ printf("FAIL at iteration %d\n", i);
+ return 1;
+ }
+ }
+
+ // CHECK: Multiple sets: PASS
+ printf("Multiple sets: PASS\n");
+
+ // Target region should execute on host (initial device)
+ int on_host = 0;
+#pragma omp target map(from : on_host)
+ {
+ on_host = omp_is_initial_device();
+ }
+
+ // CHECK: Executes on host: YES
+ printf("Executes on host: %s\n", on_host ? "YES" : "NO");
+
+ return 0;
+}
diff --git a/offload/test/offloading/target_offload_disabled_set_default_device_env.c b/offload/test/offloading/target_offload_disabled_set_default_device_env.c
new file mode 100644
index 0000000000000..3934f4211f948
--- /dev/null
+++ b/offload/test/offloading/target_offload_disabled_set_default_device_env.c
@@ -0,0 +1,45 @@
+// RUN: %libomptarget-compile-generic
+// RUN: env OMP_DEFAULT_DEVICE=5 OMP_TARGET_OFFLOAD=disabled
+// %libomptarget-run-generic 2>&1 | %fcheck-generic
+//
+// Integration test: OMP_DEFAULT_DEVICE environment variable is properly
+// overridden when offload is disabled
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ // Force runtime initialization to parse env vars
+#pragma omp parallel
+ {
+ }
+
+ int initial_device = omp_get_initial_device();
+ int default_device = omp_get_default_device();
+ int num_devices = omp_get_num_devices();
+
+ // CHECK: num_devices: 0
+ printf("num_devices: %d\n", num_devices);
+
+ // CHECK: initial_device: [[INITIAL:[0-9]+]]
+ printf("initial_device: %d\n", initial_device);
+
+ // Even though OMP_DEFAULT_DEVICE=5, should get initial_device
+ // CHECK: default_device: [[INITIAL]]
+ printf("default_device: %d\n", default_device);
+
+ // CHECK: Match: YES
+ printf("Match: %s\n", (default_device == initial_device) ? "YES" : "NO");
+
+ // Verify target region executes on host
+ int is_host = 0;
+#pragma omp target map(from : is_host)
+ {
+ is_host = omp_is_initial_device();
+ }
+
+ // CHECK: Target on host: YES
+ printf("Target on host: %s\n", is_host ? "YES" : "NO");
+
+ return 0;
+}
diff --git a/offload/test/offloading/target_offload_disabled_set_default_device_stress.c b/offload/test/offloading/target_offload_disabled_set_default_device_stress.c
new file mode 100644
index 0000000000000..b1197849b0b08
--- /dev/null
+++ b/offload/test/offloading/target_offload_disabled_set_default_device_stress.c
@@ -0,0 +1,85 @@
+// RUN: %libomptarget-compile-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled OMP_DEFAULT_DEVICE=99
+// %libomptarget-run-generic 2>&1 | %fcheck-generic
+//
+// Stress test: Multiple scenarios with offload disabled to ensure robustness
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ // Force runtime initialization
+#pragma omp parallel
+ {
+ }
+
+ int initial = omp_get_initial_device();
+
+ // CHECK: Starting test with initial_device: [[INITIAL:[0-9]+]]
+ printf("Starting test with initial_device: %d\n", initial);
+
+ // Test 1: Extreme values
+ int extreme_values[] = {-1, 0, 1, 100, 1000, 99999};
+ for (int i = 0; i < 6; i++) {
+ omp_set_default_device(extreme_values[i]);
+ int dev = omp_get_default_device();
+ if (dev != initial) {
+ printf("FAIL: extreme value %d\n", extreme_values[i]);
+ return 1;
+ }
+ }
+ // CHECK: Test 1 (extreme values): PASS
+ printf("Test 1 (extreme values): PASS\n");
+
+ // Test 2: Rapid consecutive sets
+ for (int i = 0; i < 100; i++) {
+ omp_set_default_device(i);
+ }
+ int dev = omp_get_default_device();
+ // CHECK: Test 2 (rapid sets): [[INITIAL]]
+ printf("Test 2 (rapid sets): %d\n", dev);
+
+ // Test 3: Parallel region with device operations
+ int errors = 0;
+#pragma omp parallel num_threads(4) reduction(+ : errors)
+ {
+ omp_set_default_device(omp_get_thread_num() * 50);
+ int d = omp_get_default_device();
+ if (d != initial) {
+ errors++;
+ }
+ }
+ // CHECK: Test 3 (parallel): 0 errors
+ printf("Test 3 (parallel): %d errors\n", errors);
+
+ // Test 4: Multiple target regions
+ for (int i = 0; i < 5; i++) {
+ omp_set_default_device(i * 20);
+ int executed = 0;
+#pragma omp target map(tofrom : executed)
+ {
+ executed = 1;
+ }
+ if (!executed) {
+ printf("FAIL: target %d didn't execute\n", i);
+ return 1;
+ }
+ }
+ // CHECK: Test 4 (multiple targets): PASS
+ printf("Test 4 (multiple targets): PASS\n");
+
+ // Test 5: Nested target regions
+ int outer_executed = 0;
+#pragma omp target map(tofrom : outer_executed)
+ {
+ outer_executed = 1;
+ // Note: nested target would require additional handling
+ }
+ // CHECK: Test 5 (nested targets): PASS
+ printf("Test 5 (nested targets): %s\n", outer_executed ? "PASS" : "FAIL");
+
+ // CHECK: All tests completed successfully
+ printf("All tests completed successfully\n");
+
+ return 0;
+}
diff --git a/openmp/runtime/src/kmp_ftn_entry.h b/openmp/runtime/src/kmp_ftn_entry.h
index 6adf60e7ad210..e649f876c6658 100644
--- a/openmp/runtime/src/kmp_ftn_entry.h
+++ b/openmp/runtime/src/kmp_ftn_entry.h
@@ -1143,23 +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.
-#else
- __kmp_entry_thread()->th.th_current_task->td_icvs.default_device =
- KMP_DEREF arg;
-#endif
-}
-
// Get number of NON-HOST devices.
// libomptarget, if loaded, provides this function in api.cpp.
int FTN_STDCALL KMP_EXPAND_NAME(FTN_GET_NUM_DEVICES)(void)
@@ -1197,6 +1180,31 @@ int FTN_STDCALL KMP_EXPAND_NAME(FTN_GET_INITIAL_DEVICE)(void) {
return KMP_EXPAND_NAME(FTN_GET_NUM_DEVICES)();
}
+void FTN_STDCALL KMP_EXPAND_NAME(FTN_SET_DEFAULT_DEVICE)(int KMP_DEREF arg) {
+#if KMP_MIC || KMP_OS_DARWIN || defined(KMP_STUB)
+// Nothing.
+#else
+ int device_num = KMP_DEREF arg;
+
+ // When offload is disabled, always set to initial device
+ // This prevents storing invalid device numbers in the ICV
+ if (__kmp_target_offload == tgt_disabled) {
+ device_num = KMP_EXPAND_NAME(FTN_GET_INITIAL_DEVICE)();
+ }
+
+ __kmp_entry_thread()->th.th_current_task->td_icvs.default_device = device_num;
+#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 the ICV value (which is now guaranteed to be valid by the setter)
+ 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/src/kmp_runtime.cpp b/openmp/runtime/src/kmp_runtime.cpp
index 48e29c9f9fe45..47c2364b9a195 100644
--- a/openmp/runtime/src/kmp_runtime.cpp
+++ b/openmp/runtime/src/kmp_runtime.cpp
@@ -3295,6 +3295,20 @@ static kmp_internal_control_t __kmp_get_global_icvs(void) {
KMP_DEBUG_ASSERT(__kmp_nested_proc_bind.used > 0);
+ // When offload is disabled, default_device should be the initial device
+ // regardless of OMP_DEFAULT_DEVICE environment variable or explicit set
+ int default_device_icv = __kmp_default_device;
+ if (__kmp_target_offload == tgt_disabled) {
+ // Initial device is same as num_devices when offload is disabled
+ int (*fptr)();
+ if ((*(void **)(&fptr) = KMP_DLSYM("__tgt_get_num_devices"))) {
+ default_device_icv = (*fptr)();
+ } else {
+ default_device_icv = 0;
+ }
+ }
+
+ // clang-format off
kmp_internal_control_t g_icvs = {
0, // int serial_nesting_level; //corresponds to value of th_team_serialized
(kmp_int8)__kmp_global.g.g_dynamic, // internal control for dynamic
@@ -3317,9 +3331,10 @@ static kmp_internal_control_t __kmp_get_global_icvs(void) {
r_sched, // kmp_r_sched_t sched; //internal control for runtime schedule
// {sched,chunk} pair
__kmp_nested_proc_bind.bind_types[0],
- __kmp_default_device,
+ default_device_icv,
NULL // struct kmp_internal_control *next;
};
+ // clang-format on
return g_icvs;
}
diff --git a/openmp/runtime/test/env/omp_target_offload_disabled_set_default_device.c b/openmp/runtime/test/env/omp_target_offload_disabled_set_default_device.c
new file mode 100644
index 0000000000000..629562ab43b97
--- /dev/null
+++ b/openmp/runtime/test/env/omp_target_offload_disabled_set_default_device.c
@@ -0,0 +1,72 @@
+// RUN: %libomp-compile-and-run
+//
+// Test that omp_set_default_device() stores initial_device in the ICV
+// when OMP_TARGET_OFFLOAD=DISABLED, making omp_get_default_device() return
+// a valid value.
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <omp.h>
+
+extern void kmp_set_defaults(char const *str);
+
+int main() {
+ // Disable offload first
+ kmp_set_defaults("OMP_TARGET_OFFLOAD=DISABLED");
+
+ // Force runtime initialization
+#pragma omp parallel
+ {
+ }
+
+ int initial_device = omp_get_initial_device();
+ int num_devices = omp_get_num_devices();
+
+ printf("Configuration:\n");
+ printf(" num_devices = %d\n", num_devices);
+ printf(" initial_device = %d\n", initial_device);
+
+ // Test: Set default device to invalid value
+ omp_set_default_device(5);
+
+ // Get default device - should return initial_device (not 5)
+ // because setter clamped it
+ int default_device = omp_get_default_device();
+
+ printf("After omp_set_default_device(5):\n");
+ printf(" default_device = %d\n", default_device);
+
+ if (default_device != initial_device) {
+ fprintf(stderr, "FAIL: Setter didn't clamp to initial_device\n");
+ fprintf(stderr, " Expected %d, got %d\n", initial_device,
+ default_device);
+ return EXIT_FAILURE;
+ }
+
+ // Try another value
+ omp_set_default_device(10);
+ default_device = omp_get_default_device();
+
+ printf("After omp_set_default_device(10):\n");
+ printf(" default_device = %d\n", default_device);
+
+ if (default_device != initial_device) {
+ fprintf(stderr, "FAIL: Setter didn't clamp second value\n");
+ return EXIT_FAILURE;
+ }
+
+ // Verify target region works
+ int executed = 0;
+#pragma omp target map(tofrom : executed)
+ {
+ executed = 1;
+ }
+
+ if (!executed) {
+ fprintf(stderr, "FAIL: Target region didn't execute\n");
+ return EXIT_FAILURE;
+ }
+
+ printf("PASS: Setter correctly clamps invalid device to initial_device\n");
+ return EXIT_SUCCESS;
+}
diff --git a/openmp/runtime/test/env/omp_target_offload_disabled_set_default_device_combined.c b/openmp/runtime/test/env/omp_target_offload_disabled_set_default_device_combined.c
new file mode 100644
index 0000000000000..dc8bb14ae0dfc
--- /dev/null
+++ b/openmp/runtime/test/env/omp_target_offload_disabled_set_default_device_combined.c
@@ -0,0 +1,70 @@
+// RUN: %libomp-compile-and-run
+//
+// Test combined scenario: OMP_DEFAULT_DEVICE env var + omp_set_default_device()
+// API Both should be overridden when OMP_TARGET_OFFLOAD=DISABLED
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <omp.h>
+
+extern void kmp_set_defaults(char const *str);
+
+int main() {
+ // Worst case: both env var and API try to set invalid device
+ kmp_set_defaults("OMP_DEFAULT_DEVICE=3");
+ kmp_set_defaults("OMP_TARGET_OFFLOAD=DISABLED");
+
+ // Force runtime initialization
+#pragma omp parallel
+ {
+ }
+
+ int initial_device = omp_get_initial_device();
+ int default_device = omp_get_default_device();
+
+ printf("Test 1: With OMP_DEFAULT_DEVICE=3 and offload disabled\n");
+ printf(" initial_device = %d\n", initial_device);
+ printf(" default_device = %d\n", default_device);
+
+ if (default_device != initial_device) {
+ fprintf(stderr, "FAIL: Env var not overridden\n");
+ return EXIT_FAILURE;
+ }
+
+ // Now try API call on top of env var
+ omp_set_default_device(7);
+ default_device = omp_get_default_device();
+
+ printf("\nTest 2: After additional omp_set_default_device(7)\n");
+ printf(" default_device = %d\n", default_device);
+
+ if (default_device != initial_device) {
+ fprintf(stderr, "FAIL: API call not overridden\n");
+ return EXIT_FAILURE;
+ }
+
+ // Test 3: Multiple sets in sequence
+ for (int i = 0; i < 5; i++) {
+ omp_set_default_device(i + 10);
+ default_device = omp_get_default_device();
+ if (default_device != initial_device) {
+ fprintf(stderr, "FAIL: Set %d not overridden\n", i);
+ return EXIT_FAILURE;
+ }
+ }
+
+ // Test 4: Use in target region
+ int result = -1;
+#pragma omp target device(omp_get_default_device()) map(from : result)
+ {
+ result = omp_is_initial_device();
+ }
+
+ if (result != 1) {
+ fprintf(stderr, "FAIL: Target didn't execute on initial device\n");
+ return EXIT_FAILURE;
+ }
+
+ printf("\nPASS: All combinations correctly override invalid device\n");
+ return EXIT_SUCCESS;
+}
diff --git a/openmp/runtime/test/env/omp_target_offload_disabled_set_default_device_icv.c b/openmp/runtime/test/env/omp_target_offload_disabled_set_default_device_icv.c
new file mode 100644
index 0000000000000..11c719ac105e0
--- /dev/null
+++ b/openmp/runtime/test/env/omp_target_offload_disabled_set_default_device_icv.c
@@ -0,0 +1,80 @@
+// RUN: %libomp-compile-and-run
+//
+// Test that the ICV (Internal Control Variable) itself contains the correct
+// value This verifies spec compliance: omp_get_default_device() returns the ICV
+// value, and the setter ensures the ICV is always valid.
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <omp.h>
+
+extern void kmp_set_defaults(char const *str);
+
+int main() {
+ kmp_set_defaults("OMP_TARGET_OFFLOAD=DISABLED");
+
+ // Force runtime initialization
+#pragma omp parallel
+ {
+ }
+
+ int initial_device = omp_get_initial_device();
+ printf("initial_device = %d\n", initial_device);
+
+ // Test 1: Initial value should be initial_device
+ int dev1 = omp_get_default_device();
+ printf("Test 1 - Initial: default_device = %d\n", dev1);
+ if (dev1 != initial_device) {
+ fprintf(stderr, "FAIL: Initial default != initial_device\n");
+ return EXIT_FAILURE;
+ }
+
+ // Test 2: After setting to 5, getter should still return initial_device
+ // This proves the ICV was set to initial_device (not 5)
+ omp_set_default_device(5);
+ int dev2 = omp_get_default_device();
+ printf("Test 2 - After set(5): default_device = %d\n", dev2);
+ if (dev2 != initial_device) {
+ fprintf(stderr, "FAIL: Setter didn't update ICV to initial_device\n");
+ return EXIT_FAILURE;
+ }
+
+ // Test 3: Multiple sets should all result in initial_device
+ int test_values[] = {10, 20, 100, -1, 999};
+ for (int i = 0; i < 5; i++) {
+ omp_set_default_device(test_values[i]);
+ int dev = omp_get_default_device();
+ printf("Test 3.%d - After set(%d): default_device = %d\n", i,
+ test_values[i], dev);
+ if (dev != initial_device) {
+ fprintf(stderr, "FAIL: Set(%d) resulted in %d\n", test_values[i], dev);
+ return EXIT_FAILURE;
+ }
+ }
+
+ // Test 4: Verify getter is truly just returning the ICV
+ // Call getter multiple times without any setter calls
+ for (int i = 0; i < 10; i++) {
+ int dev = omp_get_default_device();
+ if (dev != initial_device) {
+ fprintf(stderr, "FAIL: Getter call %d returned %d\n", i, dev);
+ return EXIT_FAILURE;
+ }
+ }
+
+ // Test 5: Use the device in actual operation
+ int executed = 0;
+ int device_for_target = omp_get_default_device();
+#pragma omp target device(device_for_target) map(tofrom : executed)
+ {
+ executed = 1;
+ }
+
+ if (!executed) {
+ fprintf(stderr, "FAIL: Target with default_device didn't execute\n");
+ return EXIT_FAILURE;
+ }
+
+ printf("PASS: ICV always contains valid initial_device value\n");
+ return EXIT_SUCCESS;
+}
diff --git a/openmp/runtime/test/env/omp_target_offload_disabled_set_default_device_nested.c b/openmp/runtime/test/env/omp_target_offload_disabled_set_default_device_nested.c
new file mode 100644
index 0000000000000..79851dd0e49ff
--- /dev/null
+++ b/openmp/runtime/test/env/omp_target_offload_disabled_set_default_device_nested.c
@@ -0,0 +1,87 @@
+// RUN: %libomp-compile-and-run
+//
+// Test setter behavior across nested parallel regions and 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_device(const char *context) {
+ 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;
+ }
+ return 0;
+}
+
+int main() {
+ int errors = 0;
+
+ kmp_set_defaults("OMP_DEFAULT_DEVICE=6");
+ 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
+ errors += check_device("sequential");
+
+ // Test 2: Parallel region
+#pragma omp parallel reduction(+ : errors)
+ {
+ errors += check_device("parallel");
+
+ // Each thread tries to set device
+ omp_set_default_device(omp_get_thread_num() + 20);
+ errors += check_device("after thread set");
+
+ // Test 3: Nested parallel
+#pragma omp parallel reduction(+ : errors) if (omp_get_max_threads() > 2)
+ {
+ errors += check_device("nested parallel");
+
+ // Nested thread also tries to set
+ omp_set_default_device(omp_get_thread_num() + 50);
+ errors += check_device("nested after set");
+ }
+
+#pragma omp barrier
+ errors += check_device("after barrier");
+ }
+
+ // Test 4: Back in sequential
+ errors += check_device("sequential final");
+
+ // Test 5: Target region execution (only use well-defined functions)
+ // Note: Calling omp_get_default_device() from within target region is
+ // unspecified
+ int target_errors = 0;
+#pragma omp target map(tofrom : target_errors)
+ {
+ // Verify target executes on initial device when offload disabled
+ if (!omp_is_initial_device()) {
+ target_errors = 1;
+ }
+ }
+ errors += target_errors;
+
+ if (errors > 0) {
+ fprintf(stderr, "FAIL: %d errors detected\n", errors);
+ return EXIT_FAILURE;
+ }
+
+ printf("PASS: Consistent behavior across all nesting levels\n");
+ return EXIT_SUCCESS;
+}
diff --git a/openmp/runtime/test/env/omp_target_offload_disabled_set_default_device_threads.c b/openmp/runtime/test/env/omp_target_offload_disabled_set_default_device_threads.c
new file mode 100644
index 0000000000000..e61a071628358
--- /dev/null
+++ b/openmp/runtime/test/env/omp_target_offload_disabled_set_default_device_threads.c
@@ -0,0 +1,69 @@
+// RUN: %libomp-compile-and-run
+//
+// Multi-threaded stress test: verify setter correctly clamps device numbers
+// across concurrent threads when OMP_TARGET_OFFLOAD=DISABLED
+
+#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 = 50;
+
+ kmp_set_defaults("OMP_TARGET_OFFLOAD=DISABLED");
+
+ // Force runtime initialization
+#pragma omp parallel
+ {
+ }
+
+ int initial_device = omp_get_initial_device();
+ int errors = 0;
+
+ printf("Multi-threaded test: %d threads, %d iterations each\n", NUM_THREADS,
+ NUM_ITERATIONS);
+ printf("initial_device = %d\n", initial_device);
+
+#pragma omp parallel num_threads(NUM_THREADS) reduction(+ : errors)
+ {
+ int tid = omp_get_thread_num();
+
+ for (int i = 0; i < NUM_ITERATIONS; i++) {
+ // Each thread tries to set different device
+ int attempt_device = tid * 100 + i;
+ omp_set_default_device(attempt_device);
+
+ // Should always get initial_device back
+ int default_device = omp_get_default_device();
+
+ if (default_device != initial_device) {
+#pragma omp critical
+ {
+ fprintf(stderr,
+ "FAIL: Thread %d iteration %d: set %d, got %d, expected %d\n",
+ tid, i, attempt_device, default_device, initial_device);
+ }
+ errors++;
+ }
+ }
+ }
+
+ if (errors > 0) {
+ fprintf(stderr, "FAIL: %d errors 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\n");
+ return EXIT_FAILURE;
+ }
+
+ printf("PASS: All %d threads consistently got initial_device\n",
+ NUM_THREADS * NUM_ITERATIONS);
+ return EXIT_SUCCESS;
+}
More information about the llvm-commits
mailing list