diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 09cc8c25538e..3ec09b4e8c63 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8054,11 +8054,17 @@ private:
           if (!StrideExpr)
             return false;
 
+          assert(StrideExpr->getType()->isIntegerType() &&
+                 "Stride expression must be of integer type");
+
+          // If stride is not evaluatable as a constant, treat as
+          // non-contiguous.
           const auto Constant =
               StrideExpr->getIntegerConstantExpr(CGF.getContext());
           if (!Constant)
-            return false;
+            return true;
 
+          // Treat non-unitary strides as non-contiguous.
           return !Constant->isOne();
         });
 
diff --git a/clang/test/OpenMP/target_update_variable_stride_codegen.c b/clang/test/OpenMP/target_update_variable_stride_codegen.c
new file mode 100644
index 000000000000..afc5459787c2
--- /dev/null
+++ b/clang/test/OpenMP/target_update_variable_stride_codegen.c
@@ -0,0 +1,135 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _ --global-value-regex "\.offload_.*" --global-hex-value-regex ".offload_maptypes.*" --version 3
+
+// RUN: %clang_cc1 -fopenmp -x c -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - \
+// RUN: | FileCheck %s
+
+// Check same results after serialization round-trip
+// RUN: %clang_cc1 -fopenmp -x c -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -emit-llvm %s -o - \
+// RUN: | FileCheck %s
+
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+// Test that variable stride expressions in target update directives correctly
+// set the OMP_MAP_NON_CONTIG flag. For NON_CONTIG entries, offload_sizes contains
+// dimension count (2 for 1D array section). For non-NON_CONTIG entries,
+// offload_sizes contains byte size (5 elements * 4 bytes = 20).
+int data[10];
+int stride;
+
+void test_variable_stride_to() {
+  #pragma omp target update to(data[0:5:stride])
+}
+
+void test_variable_stride_from() {
+  #pragma omp target update from(data[0:5:stride])
+}
+
+void test_constant_stride_one() {
+  #pragma omp target update to(data[0:5:1])
+}
+
+#endif // HEADER
+//.
+// CHECK: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 2]
+// CHECK: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x100000000001]]]
+// CHECK: @.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 2]
+// CHECK: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 [[#0x100000000002]]]
+// CHECK: @.offload_sizes.3 = private unnamed_addr constant [1 x i64] [i64 20]
+// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [1 x i64] [i64 [[#0x1]]]
+//.
+// CHECK-LABEL: define dso_local void @test_variable_stride_to(
+// CHECK-SAME: ) #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
+// CHECK-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8
+// CHECK-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8
+// CHECK-NEXT:    [[DIMS:%.*]] = alloca [2 x [[STRUCT_DESCRIPTOR_DIM:%.*]]], align 8
+// CHECK-NEXT:    [[TMP0:%.*]] = load i32, ptr @stride, align 4
+// CHECK-NEXT:    [[TMP1:%.*]] = zext i32 [[TMP0]] to i64
+// CHECK-NEXT:    [[TMP2:%.*]] = mul nuw i64 4, [[TMP1]]
+// CHECK-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK-NEXT:    store ptr @data, ptr [[TMP3]], align 8
+// CHECK-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-NEXT:    store ptr @data, ptr [[TMP4]], align 8
+// CHECK-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
+// CHECK-NEXT:    store ptr null, ptr [[TMP5]], align 8
+// CHECK-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [2 x [[STRUCT_DESCRIPTOR_DIM]]], ptr [[DIMS]], i64 0, i64 0
+// CHECK-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM]], ptr [[TMP6]], i32 0, i32 0
+// CHECK-NEXT:    store i64 0, ptr [[TMP7]], align 8
+// CHECK-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM]], ptr [[TMP6]], i32 0, i32 1
+// CHECK-NEXT:    store i64 5, ptr [[TMP8]], align 8
+// CHECK-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM]], ptr [[TMP6]], i32 0, i32 2
+// CHECK-NEXT:    store i64 [[TMP2]], ptr [[TMP9]], align 8
+// CHECK-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [2 x [[STRUCT_DESCRIPTOR_DIM]]], ptr [[DIMS]], i64 0, i64 1
+// CHECK-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM]], ptr [[TMP10]], i32 0, i32 0
+// CHECK-NEXT:    store i64 0, ptr [[TMP11]], align 8
+// CHECK-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM]], ptr [[TMP10]], i32 0, i32 1
+// CHECK-NEXT:    store i64 1, ptr [[TMP12]], align 8
+// CHECK-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM]], ptr [[TMP10]], i32 0, i32 2
+// CHECK-NEXT:    store i64 4, ptr [[TMP13]], align 8
+// CHECK-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-NEXT:    store ptr [[DIMS]], ptr [[TMP14]], align 8
+// CHECK-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-NEXT:    call void @__tgt_target_data_update_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 1, ptr [[TMP15]], ptr [[TMP16]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define dso_local void @test_variable_stride_from(
+// CHECK-SAME: ) #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
+// CHECK-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8
+// CHECK-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8
+// CHECK-NEXT:    [[DIMS:%.*]] = alloca [2 x [[STRUCT_DESCRIPTOR_DIM_0:%.*]]], align 8
+// CHECK-NEXT:    [[TMP0:%.*]] = load i32, ptr @stride, align 4
+// CHECK-NEXT:    [[TMP1:%.*]] = zext i32 [[TMP0]] to i64
+// CHECK-NEXT:    [[TMP2:%.*]] = mul nuw i64 4, [[TMP1]]
+// CHECK-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK-NEXT:    store ptr @data, ptr [[TMP3]], align 8
+// CHECK-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-NEXT:    store ptr @data, ptr [[TMP4]], align 8
+// CHECK-NEXT:    [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
+// CHECK-NEXT:    store ptr null, ptr [[TMP5]], align 8
+// CHECK-NEXT:    [[TMP6:%.*]] = getelementptr inbounds [2 x [[STRUCT_DESCRIPTOR_DIM_0]]], ptr [[DIMS]], i64 0, i64 0
+// CHECK-NEXT:    [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM_0]], ptr [[TMP6]], i32 0, i32 0
+// CHECK-NEXT:    store i64 0, ptr [[TMP7]], align 8
+// CHECK-NEXT:    [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM_0]], ptr [[TMP6]], i32 0, i32 1
+// CHECK-NEXT:    store i64 5, ptr [[TMP8]], align 8
+// CHECK-NEXT:    [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM_0]], ptr [[TMP6]], i32 0, i32 2
+// CHECK-NEXT:    store i64 [[TMP2]], ptr [[TMP9]], align 8
+// CHECK-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [2 x [[STRUCT_DESCRIPTOR_DIM_0]]], ptr [[DIMS]], i64 0, i64 1
+// CHECK-NEXT:    [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM_0]], ptr [[TMP10]], i32 0, i32 0
+// CHECK-NEXT:    store i64 0, ptr [[TMP11]], align 8
+// CHECK-NEXT:    [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM_0]], ptr [[TMP10]], i32 0, i32 1
+// CHECK-NEXT:    store i64 1, ptr [[TMP12]], align 8
+// CHECK-NEXT:    [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM_0]], ptr [[TMP10]], i32 0, i32 2
+// CHECK-NEXT:    store i64 4, ptr [[TMP13]], align 8
+// CHECK-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-NEXT:    store ptr [[DIMS]], ptr [[TMP14]], align 8
+// CHECK-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK-NEXT:    [[TMP16:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-NEXT:    call void @__tgt_target_data_update_mapper(ptr @[[GLOB1]], i64 -1, i32 1, ptr [[TMP15]], ptr [[TMP16]], ptr @.offload_sizes.1, ptr @.offload_maptypes.2, ptr null, ptr null)
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define dso_local void @test_constant_stride_one(
+// CHECK-SAME: ) #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
+// CHECK-NEXT:    [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8
+// CHECK-NEXT:    [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8
+// CHECK-NEXT:    [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK-NEXT:    store ptr @data, ptr [[TMP0]], align 8
+// CHECK-NEXT:    [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-NEXT:    store ptr @data, ptr [[TMP1]], align 8
+// CHECK-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
+// CHECK-NEXT:    store ptr null, ptr [[TMP2]], align 8
+// CHECK-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-NEXT:    call void @__tgt_target_data_update_mapper(ptr @[[GLOB1]], i64 -1, i32 1, ptr [[TMP3]], ptr [[TMP4]], ptr @.offload_sizes.3, ptr @.offload_maptypes.4, ptr null, ptr null)
+// CHECK-NEXT:    ret void
+//
diff --git a/offload/test/offloading/strided_update_multiple_arrays_variable_stride.c b/offload/test/offloading/strided_update_multiple_arrays_variable_stride.c
new file mode 100644
index 000000000000..32859202a20f
--- /dev/null
+++ b/offload/test/offloading/strided_update_multiple_arrays_variable_stride.c
@@ -0,0 +1,145 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Tests multiple arrays with different variable strides in single update
+// clause.
+
+#include <omp.h>
+#include <stdio.h>
+
+void test_1_update_from_multiple() {
+  int stride1 = 2;
+  int stride2 = 2;
+  double data1[10], data2[10];
+
+  // Initialize data on host
+  for (int i = 0; i < 10; i++) {
+    data1[i] = i;
+    data2[i] = i * 10;
+  }
+
+  printf("Test 1: Update FROM - Multiple arrays\n");
+
+#pragma omp target data map(to : stride1, stride2, data1[0 : 10], data2[0 : 10])
+  {
+#pragma omp target
+    {
+      for (int i = 0; i < 10; i++) {
+        data1[i] += i;
+        data2[i] += 100;
+      }
+    }
+
+#pragma omp target update from(data1[0 : 5 : stride1], data2[0 : 5 : stride2])
+  }
+
+  printf("from target data1:\n");
+  for (int i = 0; i < 10; i++)
+    printf("%f\n", data1[i]);
+
+  printf("\nfrom target data2:\n");
+  for (int i = 0; i < 10; i++)
+    printf("%f\n", data2[i]);
+}
+
+void test_2_update_to_multiple() {
+  int stride1 = 2;
+  int stride2 = 2;
+  double data1[10], data2[10];
+
+  for (int i = 0; i < 10; i++) {
+    data1[i] = i;
+    data2[i] = i * 10;
+  }
+
+  printf("\nTest 2: Update TO - Multiple arrays\n");
+
+#pragma omp target data map(tofrom : stride1, stride2, data1[0 : 10],          \
+                                data2[0 : 10])
+  {
+#pragma omp target
+    {
+      for (int i = 0; i < 10; i++) {
+        data1[i] = 100.0;
+        data2[i] = 20.0;
+      }
+    }
+
+    for (int i = 0; i < 10; i += 2) {
+      data1[i] = 10.0;
+      data2[i] = 5.0;
+    }
+
+#pragma omp target update to(data1[0 : 5 : stride1], data2[0 : 5 : stride2])
+
+#pragma omp target
+    {
+      for (int i = 0; i < 10; i++) {
+        data1[i] += 2.0;
+        data2[i] += 2.0;
+      }
+    }
+  }
+
+  printf("device data1 after update to:\n");
+  for (int i = 0; i < 10; i++)
+    printf("%f\n", data1[i]);
+
+  printf("\ndevice data2 after update to:\n");
+  for (int i = 0; i < 10; i++)
+    printf("%f\n", data2[i]);
+}
+
+int main() {
+  test_1_update_from_multiple();
+  test_2_update_to_multiple();
+  return 0;
+}
+
+// CHECK: Test 1: Update FROM - Multiple arrays
+// CHECK: from target data1:
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 1.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 3.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 5.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 16.000000
+// CHECK-NEXT: 9.000000
+
+// CHECK: from target data2:
+// CHECK-NEXT: 100.000000
+// CHECK-NEXT: 10.000000
+// CHECK-NEXT: 120.000000
+// CHECK-NEXT: 30.000000
+// CHECK-NEXT: 140.000000
+// CHECK-NEXT: 50.000000
+// CHECK-NEXT: 160.000000
+// CHECK-NEXT: 70.000000
+// CHECK-NEXT: 180.000000
+// CHECK-NEXT: 90.000000
+
+// CHECK: Test 2: Update TO - Multiple arrays
+// CHECK: device data1 after update to:
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 102.000000
+
+// CHECK: device data2 after update to:
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 22.000000
diff --git a/offload/test/offloading/strided_update_variable_count_and_stride.c b/offload/test/offloading/strided_update_variable_count_and_stride.c
new file mode 100644
index 000000000000..1e1e41653c2c
--- /dev/null
+++ b/offload/test/offloading/strided_update_variable_count_and_stride.c
@@ -0,0 +1,136 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Tests combining variable count expression AND variable stride in array
+// sections.
+
+#include <omp.h>
+#include <stdio.h>
+
+void test_1_update_from() {
+  int len = 10;
+  int stride = 2;
+  double data[len];
+
+  // Initialize data on host
+  for (int i = 0; i < len; i++) {
+    data[i] = i;
+  }
+
+  printf("Test 1: Update FROM - Variable count and stride\n");
+  printf("original values:\n");
+  for (int i = 0; i < len; i++)
+    printf("%f\n", data[i]);
+
+#pragma omp target data map(to : len, stride, data[0 : len])
+  {
+#pragma omp target
+    {
+      for (int i = 0; i < len; i++) {
+        data[i] += i;
+      }
+    }
+
+#pragma omp target update from(data[0 : len / 2 : stride])
+  }
+
+  printf("from target results:\n");
+  for (int i = 0; i < len; i++)
+    printf("%f\n", data[i]);
+}
+
+void test_2_update_to() {
+  int len = 10;
+  int stride = 2;
+  double data[len];
+
+  for (int i = 0; i < len; i++) {
+    data[i] = i;
+  }
+
+  printf("\nTest 2: Update TO - Variable count and stride\n");
+  printf("original values:\n");
+  for (int i = 0; i < len; i++)
+    printf("%f\n", data[i]);
+
+#pragma omp target data map(tofrom : len, stride, data[0 : len])
+  {
+#pragma omp target
+    {
+      for (int i = 0; i < len; i++) {
+        data[i] = 50.0;
+      }
+    }
+
+    for (int i = 0; i < len / 2; i++) {
+      data[i * stride] = 10.0;
+    }
+
+#pragma omp target update to(data[0 : len / 2 : stride])
+
+#pragma omp target
+    {
+      for (int i = 0; i < len; i++) {
+        data[i] += 5.0;
+      }
+    }
+  }
+
+  printf("device values after update to:\n");
+  for (int i = 0; i < len; i++)
+    printf("%f\n", data[i]);
+}
+
+int main() {
+  test_1_update_from();
+  test_2_update_to();
+  return 0;
+}
+
+// CHECK: Test 1: Update FROM - Variable count and stride
+// CHECK: original values:
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 1.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 3.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 5.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 9.000000
+
+// CHECK: from target results:
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 1.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 3.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 5.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 16.000000
+// CHECK-NEXT: 9.000000
+
+// CHECK: Test 2: Update TO - Variable count and stride
+// CHECK: original values:
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 1.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 3.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 5.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 9.000000
+
+// CHECK: device values after update to:
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
diff --git a/offload/test/offloading/strided_update_variable_stride.c b/offload/test/offloading/strided_update_variable_stride.c
new file mode 100644
index 000000000000..7c8079efa56e
--- /dev/null
+++ b/offload/test/offloading/strided_update_variable_stride.c
@@ -0,0 +1,134 @@
+// This test checks "update from" and "update to" with variable stride.
+// Tests data[0:5:stride] where stride is a variable, making it non-contiguous.
+
+// RUN: %libomptarget-compile-run-and-check-generic
+#include <omp.h>
+#include <stdio.h>
+
+void test_1_update_from() {
+  int stride = 2;
+  double data[10];
+
+  // Initialize data on host
+  for (int i = 0; i < 10; i++) {
+    data[i] = i;
+  }
+
+  printf("Test 1: Update FROM device\n");
+  printf("original values:\n");
+  for (int i = 0; i < 10; i++)
+    printf("%f\n", data[i]);
+
+#pragma omp target data map(to : stride, data[0 : 10])
+  {
+#pragma omp target
+    {
+      for (int i = 0; i < 10; i++) {
+        data[i] += i;
+      }
+    }
+
+#pragma omp target update from(data[0 : 5 : stride])
+  }
+
+  printf("from target results:\n");
+  for (int i = 0; i < 10; i++)
+    printf("%f\n", data[i]);
+}
+
+void test_2_update_to() {
+  int stride = 2;
+  double data[10];
+
+  for (int i = 0; i < 10; i++) {
+    data[i] = i;
+  }
+
+  printf("\nTest 2: Update TO device\n");
+  printf("original values:\n");
+  for (int i = 0; i < 10; i++)
+    printf("%f\n", data[i]);
+
+#pragma omp target data map(tofrom : stride, data[0 : 10])
+  {
+#pragma omp target
+    {
+      for (int i = 0; i < 10; i++) {
+        data[i] = 50.0;
+      }
+    }
+
+    for (int i = 0; i < 10; i += 2) {
+      data[i] = 10.0;
+    }
+
+#pragma omp target update to(data[0 : 5 : stride])
+
+#pragma omp target
+    {
+      for (int i = 0; i < 10; i++) {
+        data[i] += 5.0;
+      }
+    }
+  }
+
+  printf("device values after update to:\n");
+  for (int i = 0; i < 10; i++)
+    printf("%f\n", data[i]);
+}
+
+int main() {
+  test_1_update_from();
+  test_2_update_to();
+  return 0;
+}
+
+// CHECK: Test 1: Update FROM device
+// CHECK: original values:
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 1.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 3.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 5.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 9.000000
+
+// CHECK: from target results:
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 1.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 3.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 5.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 16.000000
+// CHECK-NEXT: 9.000000
+
+// CHECK: Test 2: Update TO device
+// CHECK: original values:
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 1.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 3.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 5.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 9.000000
+
+// CHECK: device values after update to:
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
diff --git a/offload/test/offloading/strided_update_variable_stride_misc.c b/offload/test/offloading/strided_update_variable_stride_misc.c
new file mode 100644
index 000000000000..df2bfa64cfe3
--- /dev/null
+++ b/offload/test/offloading/strided_update_variable_stride_misc.c
@@ -0,0 +1,88 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Miscellaneous variable stride tests: stride=1, stride=array_size, stride from
+// array subscript.
+
+#include <omp.h>
+#include <stdio.h>
+
+void test_1_variable_stride_one() {
+  int stride_one = 1;
+  double data1[10];
+
+  // Initialize data on host
+  for (int i = 0; i < 10; i++) {
+    data1[i] = i;
+  }
+
+#pragma omp target data map(to : stride_one, data1[0 : 10])
+  {
+#pragma omp target
+    {
+      for (int i = 0; i < 10; i++) {
+        data1[i] += i;
+      }
+    }
+
+#pragma omp target update from(data1[0 : 10 : stride_one])
+  }
+
+  printf("Test 1: Variable stride = 1\n");
+  for (int i = 0; i < 10; i++)
+    printf("%f\n", data1[i]);
+}
+
+void test_2_variable_stride_large() {
+  int stride_large = 5;
+  double data2[10];
+
+  // Initialize data on host
+  for (int i = 0; i < 10; i++) {
+    data2[i] = i;
+  }
+
+#pragma omp target data map(to : stride_large, data2[0 : 10])
+  {
+#pragma omp target
+    {
+      for (int i = 0; i < 10; i++) {
+        data2[i] += i;
+      }
+    }
+
+#pragma omp target update from(data2[0 : 2 : stride_large])
+  }
+
+  printf("\nTest 2: Variable stride = 5\n");
+  for (int i = 0; i < 10; i++)
+    printf("%f\n", data2[i]);
+}
+
+int main() {
+  test_1_variable_stride_one();
+  test_2_variable_stride_large();
+  return 0;
+}
+
+// CHECK: Test 1: Variable stride = 1
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 1.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 3.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 5.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 9.000000
+
+// CHECK: Test 2: Variable stride = 5
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 1.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 3.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 10.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 9.000000
diff --git a/offload/test/offloading/target_update_ptr_variable_count_and_stride.c b/offload/test/offloading/target_update_ptr_variable_count_and_stride.c
new file mode 100644
index 000000000000..1a28595969c6
--- /dev/null
+++ b/offload/test/offloading/target_update_ptr_variable_count_and_stride.c
@@ -0,0 +1,94 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Tests heap-allocated pointers with both variable count expression and
+// variable stride.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+int main() {
+  int len = 10;
+  int stride = 2;
+  double *result = (double *)malloc(len * sizeof(double));
+
+  for (int i = 0; i < len; i++) {
+    result[i] = 0;
+  }
+
+#pragma omp target enter data map(to : len, stride, result[0 : len])
+
+#pragma omp target map(alloc : result[0 : len])
+  {
+    for (int i = 0; i < len; i++) {
+      result[i] = i;
+    }
+  }
+
+  // Test FROM: Variable count and stride
+#pragma omp target map(alloc : result[0 : len])
+  {
+    for (int i = 0; i < len; i++) {
+      result[i] += i * 10;
+    }
+  }
+
+#pragma omp target update from(result[0 : len / 2 : stride])
+
+  printf("heap ptr variable count and stride (from):\n");
+  for (int i = 0; i < len; i++)
+    printf("%f\n", result[i]);
+
+  // Test TO: Reset, modify host, update to device
+#pragma omp target map(alloc : result[0 : len])
+  {
+    for (int i = 0; i < len; i++) {
+      result[i] = i * 2;
+    }
+  }
+
+  for (int i = 0; i < len / 2; i++) {
+    result[i * stride] = i + 100;
+  }
+
+#pragma omp target update to(result[0 : len / 2 : stride])
+
+#pragma omp target map(alloc : result[0 : len])
+  {
+    for (int i = 0; i < len; i++) {
+      result[i] += 100;
+    }
+  }
+
+#pragma omp target update from(result[0 : len])
+
+  printf("heap ptr variable count and stride (to):\n");
+  for (int i = 0; i < len; i++)
+    printf("%f\n", result[i]);
+
+#pragma omp target exit data map(delete : len, stride, result[0 : len])
+  free(result);
+  return 0;
+}
+
+// CHECK: heap ptr variable count and stride (from):
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 44.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 66.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 88.000000
+// CHECK-NEXT: 0.000000
+// CHECK: heap ptr variable count and stride (to):
+// CHECK-NEXT: 200.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 201.000000
+// CHECK-NEXT: 106.000000
+// CHECK-NEXT: 202.000000
+// CHECK-NEXT: 110.000000
+// CHECK-NEXT: 203.000000
+// CHECK-NEXT: 114.000000
+// CHECK-NEXT: 204.000000
+// CHECK-NEXT: 118.000000
diff --git a/offload/test/offloading/target_update_ptr_variable_stride.c b/offload/test/offloading/target_update_ptr_variable_stride.c
new file mode 100644
index 000000000000..bea396065b76
--- /dev/null
+++ b/offload/test/offloading/target_update_ptr_variable_stride.c
@@ -0,0 +1,95 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Tests non-contiguous array sections with variable stride on heap-allocated
+// pointers.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+int main() {
+  int stride = 2;
+  int len = 10;
+  double *result = (double *)malloc(len * sizeof(double));
+
+  // Initialize
+  for (int i = 0; i < len; i++) {
+    result[i] = 0;
+  }
+
+#pragma omp target enter data map(to : stride, len, result[0 : len])
+
+#pragma omp target map(alloc : result[0 : len])
+  {
+    for (int i = 0; i < len; i++) {
+      result[i] = i;
+    }
+  }
+
+  // Test FROM
+#pragma omp target map(alloc : result[0 : len])
+  {
+    for (int i = 0; i < len; i++) {
+      result[i] += i * 10;
+    }
+  }
+
+#pragma omp target update from(result[0 : 5 : stride])
+
+  printf("heap ptr variable stride (from):\n");
+  for (int i = 0; i < len; i++)
+    printf("%f\n", result[i]);
+
+  // Test TO: Reset, modify host, update to device
+#pragma omp target map(alloc : result[0 : len])
+  {
+    for (int i = 0; i < len; i++) {
+      result[i] = i * 2;
+    }
+  }
+
+  for (int i = 0; i < 5; i++) {
+    result[i * stride] = i + 100;
+  }
+
+#pragma omp target update to(result[0 : 5 : stride])
+
+#pragma omp target map(alloc : result[0 : len])
+  {
+    for (int i = 0; i < len; i++) {
+      result[i] += 100;
+    }
+  }
+
+#pragma omp target update from(result[0 : len])
+
+  printf("heap ptr variable stride (to):\n");
+  for (int i = 0; i < len; i++)
+    printf("%f\n", result[i]);
+
+#pragma omp target exit data map(delete : stride, len, result[0 : len])
+  free(result);
+  return 0;
+}
+
+// CHECK: heap ptr variable stride (from):
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 44.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 66.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 88.000000
+// CHECK-NEXT: 0.000000
+// CHECK: heap ptr variable stride (to):
+// CHECK-NEXT: 200.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 201.000000
+// CHECK-NEXT: 106.000000
+// CHECK-NEXT: 202.000000
+// CHECK-NEXT: 110.000000
+// CHECK-NEXT: 203.000000
+// CHECK-NEXT: 114.000000
+// CHECK-NEXT: 204.000000
+// CHECK-NEXT: 118.000000
diff --git a/offload/test/offloading/target_update_strided_struct_variable_count_and_stride.c b/offload/test/offloading/target_update_strided_struct_variable_count_and_stride.c
new file mode 100644
index 000000000000..6daf10383e92
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_variable_count_and_stride.c
@@ -0,0 +1,96 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Tests struct member arrays with both variable count expression and variable
+// stride.
+
+#include <omp.h>
+#include <stdio.h>
+
+struct S {
+  int len;
+  int stride;
+  double data[20];
+};
+
+int main() {
+  struct S s;
+  s.len = 10;
+  s.stride = 2;
+
+  // Initialize
+#pragma omp target map(tofrom : s)
+  {
+    for (int i = 0; i < s.len; i++) {
+      s.data[i] = i;
+    }
+  }
+
+  // Test FROM: Variable count and stride
+#pragma omp target data map(to : s)
+  {
+#pragma omp target
+    {
+      for (int i = 0; i < s.len; i++) {
+        s.data[i] += i * 10;
+      }
+    }
+
+#pragma omp target update from(s.data[0 : s.len / 2 : s.stride])
+  }
+
+  printf("struct variable count and stride (from):\n");
+  for (int i = 0; i < s.len; i++)
+    printf("%f\n", s.data[i]);
+
+  // Test TO: Reset, modify host, update to device
+#pragma omp target map(tofrom : s)
+  {
+    for (int i = 0; i < s.len; i++) {
+      s.data[i] = i * 2;
+    }
+  }
+
+  for (int i = 0; i < s.len / 2; i++) {
+    s.data[i * s.stride] = i + 100;
+  }
+
+#pragma omp target data map(to : s)
+  {
+#pragma omp target update to(s.data[0 : s.len / 2 : s.stride])
+
+#pragma omp target
+    {
+      for (int i = 0; i < s.len; i++) {
+        s.data[i] += 100;
+      }
+    }
+  }
+
+  printf("struct variable count and stride (to):\n");
+  for (int i = 0; i < s.len; i++)
+    printf("%f\n", s.data[i]);
+
+  return 0;
+}
+
+// CHECK: struct variable count and stride (from):
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 11.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 33.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 77.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 9.000000
+// CHECK: struct variable count and stride (to):
+// CHECK-NEXT: 100.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 101.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 10.000000
+// CHECK-NEXT: 103.000000
+// CHECK-NEXT: 14.000000
+// CHECK-NEXT: 104.000000
+// CHECK-NEXT: 18.000000
diff --git a/offload/test/offloading/target_update_strided_struct_variable_stride.c b/offload/test/offloading/target_update_strided_struct_variable_stride.c
new file mode 100644
index 000000000000..4cd9da629ca9
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_variable_stride.c
@@ -0,0 +1,95 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Tests non-contiguous array sections with variable stride on struct member
+// arrays.
+
+#include <omp.h>
+#include <stdio.h>
+
+struct S {
+  int stride;
+  double data[20];
+};
+
+int main() {
+  struct S s;
+  s.stride = 2;
+  int len = 10;
+
+  // Initialize
+#pragma omp target map(tofrom : s, len)
+  {
+    for (int i = 0; i < len; i++) {
+      s.data[i] = i;
+    }
+  }
+
+  // Test FROM
+#pragma omp target data map(to : s, len)
+  {
+#pragma omp target
+    {
+      for (int i = 0; i < len; i++) {
+        s.data[i] += i * 10;
+      }
+    }
+
+#pragma omp target update from(s.data[0 : 5 : s.stride])
+  }
+
+  printf("struct variable stride (from):\n");
+  for (int i = 0; i < len; i++)
+    printf("%f\n", s.data[i]);
+
+  // Test TO: Reset, modify host, update to device
+#pragma omp target map(tofrom : s)
+  {
+    for (int i = 0; i < len; i++) {
+      s.data[i] = i * 2;
+    }
+  }
+
+  for (int i = 0; i < 5; i++) {
+    s.data[i * s.stride] = i + 100;
+  }
+
+#pragma omp target data map(to : s)
+  {
+#pragma omp target update to(s.data[0 : 5 : s.stride])
+
+#pragma omp target
+    {
+      for (int i = 0; i < len; i++) {
+        s.data[i] += 100;
+      }
+    }
+  }
+
+  printf("struct variable stride (to):\n");
+  for (int i = 0; i < len; i++)
+    printf("%f\n", s.data[i]);
+
+  return 0;
+}
+
+// CHECK: struct variable stride (from):
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 11.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 33.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 77.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 9.000000
+// CHECK: struct variable stride (to):
+// CHECK-NEXT: 100.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 101.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 10.000000
+// CHECK-NEXT: 103.000000
+// CHECK-NEXT: 14.000000
+// CHECK-NEXT: 104.000000
+// CHECK-NEXT: 18.000000