[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
Fri Dec 26 03:04:52 PST 2025
https://github.com/amitamd7 created https://github.com/llvm/llvm-project/pull/173630
Fix a crash in OpenMP programs when `OMP_TARGET_OFFLOAD=DISABLED` is set together with a non-zero `OMP_DEFAULT_DEVICE` value or when `omp_set_default_device()` is called with an invalid device number.
**Issue:**
When users set both:
`OMP_TARGET_OFFLOAD=DISABLED` (to disable GPU offloading)
`OMP_DEFAULT_DEVICE=N` (where `N > 0`)
Or call `omp_set_default_device(N)` where `N` is invalid, the application would crash with:
`omptarget fatal error 2: "invalid value" device number 'N' out of range, only 0 devices available`
**Root Cause:**
The default-device-var Internal Control Variable (ICV) could be set to an invalid device number, which would later cause crashes when the value was used in target constructs or device queries.
This patch implements a setter-based fix that ensures the default-device-var ICV always contains a valid value when offloading is disabled.
Testing: Added testcases that trigger the issue and confirm the approach resolves the issue..
>From 11d6a80cef539af80505e09da96fdef07a960c8a 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 | 51 +++++++++++
..._offload_disabled_set_default_device_api.c | 64 ++++++++++++++
..._offload_disabled_set_default_device_env.c | 44 ++++++++++
...fload_disabled_set_default_device_stress.c | 84 ++++++++++++++++++
openmp/runtime/src/kmp_ftn_entry.h | 42 +++++----
openmp/runtime/src/kmp_runtime.cpp | 53 +++++++-----
...rget_offload_disabled_set_default_device.c | 71 +++++++++++++++
...oad_disabled_set_default_device_combined.c | 70 +++++++++++++++
..._offload_disabled_set_default_device_icv.c | 80 +++++++++++++++++
...fload_disabled_set_default_device_nested.c | 86 +++++++++++++++++++
...load_disabled_set_default_device_threads.c | 69 +++++++++++++++
11 files changed, 677 insertions(+), 37 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..5a44b04373d51
--- /dev/null
+++ b/offload/test/api/omp_set_default_device_disabled.c
@@ -0,0 +1,51 @@
+// 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..6a6beac03cd13
--- /dev/null
+++ b/offload/test/offloading/target_offload_disabled_set_default_device_api.c
@@ -0,0 +1,64 @@
+// 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..50949b6d98213
--- /dev/null
+++ b/offload/test/offloading/target_offload_disabled_set_default_device_env.c
@@ -0,0 +1,44 @@
+// 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..d3b24436d4e77
--- /dev/null
+++ b/offload/test/offloading/target_offload_disabled_set_default_device_stress.c
@@ -0,0 +1,84 @@
+// 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..0bbe13508cbf7 100644
--- a/openmp/runtime/src/kmp_runtime.cpp
+++ b/openmp/runtime/src/kmp_runtime.cpp
@@ -3295,30 +3295,43 @@ 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;
+ }
+ }
+
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
- // adjustment of threads (per thread)
- (kmp_int8)__kmp_env_blocktime, // int bt_set; //internal control for
- // whether blocktime is explicitly set
- __kmp_dflt_blocktime, // int blocktime; //internal control for blocktime
+ 0, // int serial_nesting_level; //corresponds to value of
+ // th_team_serialized
+ (kmp_int8)__kmp_global.g.g_dynamic, // internal control for dynamic
+ // adjustment of threads (per thread)
+ (kmp_int8)__kmp_env_blocktime, // int bt_set; //internal control for
+ // whether blocktime is explicitly set
+ __kmp_dflt_blocktime, // int blocktime; //internal control for blocktime
#if KMP_USE_MONITOR
- __kmp_bt_intervals, // int bt_intervals; //internal control for blocktime
+ __kmp_bt_intervals, // int bt_intervals; //internal control for blocktime
// intervals
#endif
- __kmp_dflt_team_nth, // int nproc; //internal control for # of threads for
- // next parallel region (per thread)
- // (use a max ub on value if __kmp_parallel_initialize not called yet)
- __kmp_cg_max_nth, // int thread_limit;
- __kmp_task_max_nth, // int task_thread_limit; // to set the thread_limit
- // on task. This is used in the case of target thread_limit
- __kmp_dflt_max_active_levels, // int max_active_levels; //internal control
- // for max_active_levels
- 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,
- NULL // struct kmp_internal_control *next;
+ __kmp_dflt_team_nth, // int nproc; //internal control for # of threads for
+ // next parallel region (per thread)
+ // (use a max ub on value if __kmp_parallel_initialize not called yet)
+ __kmp_cg_max_nth, // int thread_limit;
+ __kmp_task_max_nth, // int task_thread_limit; // to set the thread_limit
+ // on task. This is used in the case of target thread_limit
+ __kmp_dflt_max_active_levels, // int max_active_levels; //internal control
+ // for max_active_levels
+ r_sched, // kmp_r_sched_t sched; //internal control for runtime schedule
+ // {sched,chunk} pair
+ __kmp_nested_proc_bind.bind_types[0], default_device_icv,
+ NULL // struct kmp_internal_control *next;
};
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..253278cb62920
--- /dev/null
+++ b/openmp/runtime/test/env/omp_target_offload_disabled_set_default_device.c
@@ -0,0 +1,71 @@
+// 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..8a9196fa68ac9
--- /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..e5793ee7310f4
--- /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..cfcaab49847ba
--- /dev/null
+++ b/openmp/runtime/test/env/omp_target_offload_disabled_set_default_device_nested.c
@@ -0,0 +1,86 @@
+// 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..b84c08986bc94
--- /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