diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp
index 49f23bccbc68d..fe845d5a456ba 100644
--- a/clang/lib/CodeGen/CGCall.cpp
+++ b/clang/lib/CodeGen/CGCall.cpp
@@ -2850,8 +2850,10 @@ void CodeGenModule::ConstructAttributeList(StringRef Name,
     // > For arguments to a __kernel function declared to be a pointer to a
     // > data type, the OpenCL compiler can assume that the pointee is always
     // > appropriately aligned as required by the data type.
-    if (TargetDecl && TargetDecl->hasAttr<OpenCLKernelAttr>() &&
-        ParamType->isPointerType()) {
+    //
+    // Don't do this for SYCL, as this assumption does not hold.
+    if (!getLangOpts().SYCLIsDevice && TargetDecl &&
+        TargetDecl->hasAttr<OpenCLKernelAttr>() && ParamType->isPointerType()) {
       QualType PTy = ParamType->getPointeeType();
       if (!PTy->isIncompleteType() && PTy->isConstantSizeType()) {
         llvm::Align Alignment =
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index fc5abb02d75f3..b34d51c83a5d7 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -670,5 +670,8 @@ class image {
   }
 };
 
+template <typename T> T *malloc_shared(unsigned long size);
+void free(void *ptr);
+
 } // namespace _V1
 } // namespace sycl
diff --git a/clang/test/CodeGenSYCL/accessor-readonly.cpp b/clang/test/CodeGenSYCL/accessor-readonly.cpp
index 235d2f71a5bf7..1a2c1950a6828 100644
--- a/clang/test/CodeGenSYCL/accessor-readonly.cpp
+++ b/clang/test/CodeGenSYCL/accessor-readonly.cpp
@@ -15,7 +15,7 @@ void f0(sycl::queue &myQueue, sycl::buffer<int, 1> &in_buf, sycl::buffer<int, 1>
 // CHECK: spir_kernel{{.*}}f1_kernel
 // CHECK-NOT: readonly
 // CHECK-SAME: %_arg_write_acc{{.*}}%_arg_write_acc1{{.*}}%_arg_write_acc2{{.*}}%_arg_write_acc3
-// CHECK-SAME:  readonly align 4 %_arg_read_acc
+// CHECK-SAME:  readonly %_arg_read_acc
 void f1(sycl::queue &myQueue, sycl::buffer<int, 1> &in_buf, sycl::buffer<int, 1> &out_buf) {
   myQueue.submit([&](sycl::handler &cgh) {
     auto write_acc = out_buf.get_access<sycl::access::mode::write>(cgh);
@@ -25,7 +25,7 @@ void f1(sycl::queue &myQueue, sycl::buffer<int, 1> &in_buf, sycl::buffer<int, 1>
 }
 
 // CHECK: spir_kernel{{.*}}f2_kernel
-// CHECK-SAME: readonly align 4 %_arg_read_acc
+// CHECK-SAME: readonly %_arg_read_acc
 // CHECK-NOT: readonly
 // CHECK-SAME: %_arg_write_acc
 void f2(sycl::queue &myQueue, sycl::buffer<int, 1> &in_buf, sycl::buffer<int, 1> &out_buf) {
diff --git a/clang/test/CodeGenSYCL/accessor_inheritance.cpp b/clang/test/CodeGenSYCL/accessor_inheritance.cpp
index e8bc2711435d9..d2e5c568c07f3 100644
--- a/clang/test/CodeGenSYCL/accessor_inheritance.cpp
+++ b/clang/test/CodeGenSYCL/accessor_inheritance.cpp
@@ -26,11 +26,11 @@ int main() {
 // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE6kernel
 // CHECK-SAME: i32 noundef [[ARG_A:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: i32 noundef [[ARG_B:%[a-zA-Z0-9_]+]],
-// CHECK-SAME: ptr addrspace(1) noundef readonly align 1 [[ACC1_DATA:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr addrspace(1) noundef readonly [[ACC1_DATA:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC1_RANGE1:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC1_RANGE2:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval(%[[ID_TYPE]]) align 4 [[ACC1_ID:%[a-zA-Z0-9_]+]],
-// CHECK-SAME: ptr addrspace(1) noundef readonly align 1 [[ACC2_DATA:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr addrspace(1) noundef readonly [[ACC2_DATA:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC2_RANGE1:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC2_RANGE2:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval(%[[ID_TYPE]]) align 4 [[ACC2_ID:%[a-zA-Z0-9_]+]],
diff --git a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
index 1b296ce5200a8..66ebd2aaf269c 100644
--- a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
+++ b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
@@ -20,7 +20,7 @@ int main() {
 }
 
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_function
-// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE:%[a-zA-Z0-9_]+1]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE:%[a-zA-Z0-9_]+2]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET:%[a-zA-Z0-9_]+]])
diff --git a/clang/test/CodeGenSYCL/intel-restrict.cpp b/clang/test/CodeGenSYCL/intel-restrict.cpp
index 9845da905e0e2..65cc838f92f31 100644
--- a/clang/test/CodeGenSYCL/intel-restrict.cpp
+++ b/clang/test/CodeGenSYCL/intel-restrict.cpp
@@ -11,7 +11,7 @@ int main() {
   int *c;
   kernel<class kernel_restrict>(
       [ a, b, c ]() [[intel::kernel_args_restrict]] { c[0] = a[0] + b[0]; });
-  // CHECK: define {{.*}}spir_kernel {{.*}}kernel_restrict(ptr addrspace(1) noalias noundef align 4 %{{.*}}, ptr addrspace(1) noalias noundef align 4 %{{.*}}, ptr addrspace(1) noalias noundef align 4 %{{.*}})
+  // CHECK: define {{.*}}spir_kernel {{.*}}kernel_restrict(ptr addrspace(1) noalias noundef %{{.*}}, ptr addrspace(1) noalias noundef %{{.*}}, ptr addrspace(1) noalias noundef %{{.*}})
 
   int *d;
   int *e;
@@ -19,10 +19,10 @@ int main() {
 
   kernel<class kernel_norestrict>(
       [d, e, f]() { f[0] = d[0] + e[0]; });
-  // CHECK: define {{.*}}spir_kernel {{.*}}kernel_norestrict(ptr addrspace(1) noundef align 4 %{{.*}}, ptr addrspace(1) noundef align 4 %{{.*}}, ptr addrspace(1) noundef align 4 %{{.*}})
+  // CHECK: define {{.*}}spir_kernel {{.*}}kernel_norestrict(ptr addrspace(1) noundef %{{.*}}, ptr addrspace(1) noundef %{{.*}}, ptr addrspace(1) noundef %{{.*}})
 
   int g = 42;
   kernel<class kernel_restrict_other_types>(
       [ a, b, c, g ]() [[intel::kernel_args_restrict]] { c[0] = a[0] + b[0] + g; });
-  // CHECK: define {{.*}}spir_kernel {{.*}}kernel_restrict_other_types(ptr addrspace(1) noalias noundef align 4 %{{.*}}, ptr addrspace(1) noalias noundef align 4 %{{.*}}, ptr addrspace(1) noalias noundef align 4 %{{.*}}, i32 noundef %{{.*}})
+  // CHECK: define {{.*}}spir_kernel {{.*}}kernel_restrict_other_types(ptr addrspace(1) noalias noundef %{{.*}}, ptr addrspace(1) noalias noundef %{{.*}}, ptr addrspace(1) noalias noundef %{{.*}}, i32 noundef %{{.*}})
 }
diff --git a/clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp b/clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp
index db432d97e5b94..b72da1488d477 100644
--- a/clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp
+++ b/clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp
@@ -102,11 +102,11 @@ int main() {
 
 // Check kernel_A parameters
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A
-// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG1:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]],
-// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG2:%[a-zA-Z0-9_]+4]],
+// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG2:%[a-zA-Z0-9_]+4]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+6]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+7]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+8]])
@@ -115,7 +115,7 @@ int main() {
 
 // Check kernel_readOnlyAcc parameters
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_readOnlyAcc
-// CHECK-SAME: ptr addrspace(1) noundef readonly align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr addrspace(1) noundef readonly [[MEM_ARG1:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
@@ -135,13 +135,13 @@ int main() {
 
 // Check usm_ptr parameters
 // CHECK: define {{.*}}spir_kernel void @{{.*}}usm_ptr
-// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
-// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]]
+// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG1:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG1:%[a-zA-Z0-9_]+]]
 // CHECK-NOT: kernel_arg_runtime_aligned
 // CHECK-NOT: kernel_arg_exclusive_ptr
 
 // CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessorDep
-// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG1:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
@@ -149,7 +149,7 @@ int main() {
 // CHECK-SAME: !kernel_arg_exclusive_ptr ![[#ACCESSORMD2]]
 
 // CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessor
-// CHECK-SAME: ptr addrspace(3) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr addrspace(3) noundef [[MEM_ARG1:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
@@ -158,11 +158,11 @@ int main() {
 
 // Check kernel_acc_raw_ptr parameters
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_acc_raw_ptr
-// CHECK-SAME: ptr addrspace(1) noundef readonly align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr addrspace(1) noundef readonly [[MEM_ARG1:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
-// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]]
+// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG1:%[a-zA-Z0-9_]+]]
 // CHECK-SAME: !kernel_arg_runtime_aligned ![[#ACCESSORMD3:]]
 // CHECK-SAME: !kernel_arg_exclusive_ptr ![[#ACCESSORMD3]]
 
diff --git a/clang/test/CodeGenSYCL/kernel-arg-align.cpp b/clang/test/CodeGenSYCL/kernel-arg-align.cpp
new file mode 100644
index 0000000000000..3a55e6e5c7b3c
--- /dev/null
+++ b/clang/test/CodeGenSYCL/kernel-arg-align.cpp
@@ -0,0 +1,52 @@
+// RUN: %clang_cc1 -fsycl-is-device -O0 -internal-isystem %S/Inputs -triple spir64 -emit-llvm -o - %s | FileCheck %s
+
+// Test that the pointer parameters generated for the kernel do not
+// have alignment on them.
+
+#include "sycl.hpp"
+
+using namespace sycl;
+
+struct S;
+
+void Test() {
+  struct MyIP {
+    char* a;
+    int* b;
+    double* c;
+
+    void operator()() const {
+       *((int *) a)  = 1; // 1 on arg, 4 on site
+       *((double *) b)  = 2; // 4 on arg, 8 on site
+       *((char *) c)  = 3; // 8 on arg, 1 on site
+    }
+  };
+
+  constexpr int kN = 8;
+  auto host_array_A =
+      malloc_shared<char>(kN);
+
+  auto host_array_B =
+      malloc_shared<int>(kN);
+
+  auto host_array_C =
+      malloc_shared<double>(kN);
+
+  for (int i = 0; i < kN; i++) {
+    host_array_A[i] = i;
+    host_array_B[i] = i * 2;
+  }
+
+  sycl::kernel_single_task<S>(MyIP{host_array_A, host_array_B, host_array_C});
+
+  free(host_array_A);
+  free(host_array_B);
+  free(host_array_C);
+}
+
+int main() {
+  Test();
+  return 0;
+}
+
+// CHECK: define {{.*}} spir_kernel void @_ZTS1S(ptr addrspace(1) noundef %_arg_a, ptr addrspace(1) noundef %_arg_b, ptr addrspace(1) noundef %_arg_c)
diff --git a/clang/test/CodeGenSYCL/kernel-handler.cpp b/clang/test/CodeGenSYCL/kernel-handler.cpp
index efa27788fd57c..a738c3d1e1d98 100644
--- a/clang/test/CodeGenSYCL/kernel-handler.cpp
+++ b/clang/test/CodeGenSYCL/kernel-handler.cpp
@@ -23,7 +23,7 @@ void test(int val) {
 }
 
 // ALL: define dso_local{{ spir_kernel | }}void @{{.*}}test_kernel_handler{{[^(]*}}
-// ALL-SAME: (i32 noundef %_arg_a, ptr addrspace(1) noundef align 1 %_arg__specialization_constants_buffer)
+// ALL-SAME: (i32 noundef %_arg_a, ptr addrspace(1) noundef %_arg__specialization_constants_buffer)
 // ALL: %kh = alloca %"class.sycl::_V1::kernel_handler", align 1
 
 // NONATIVESUPPORT: %[[KH:[0-9]+]] = load ptr addrspace(1), ptr %_arg__specialization_constants_buffer.addr, align 8
diff --git a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp
index e2ed6126f9bab..ab3fd03f90944 100644
--- a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp
@@ -25,11 +25,11 @@ int main() {
 
 // Check kernel_A parameters
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A
-// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG1:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]],
-// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG2:%[a-zA-Z0-9_]+4]],
+// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG2:%[a-zA-Z0-9_]+4]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+6]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+7]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+8]])
diff --git a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp
index b569e152899f6..f1ca658fb9abe 100644
--- a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp
@@ -28,11 +28,11 @@ int main() {
 
 // CHECK kernel_C parameters
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_C
-// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG1:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval({{.*}}) align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
 // CHECK-SAME: ptr noundef byval({{.*}}) align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
 // CHECK-SAME: ptr noundef byval({{.*}}) align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]],
-// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG2:%[a-zA-Z0-9_]+4]],
+// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG2:%[a-zA-Z0-9_]+4]],
 // CHECK-SAME: ptr noundef byval({{.*}}) align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+6]],
 // CHECK-SAME: ptr noundef byval({{.*}}) align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+7]],
 // CHECK-SAME: ptr noundef byval({{.*}}) align 4 [[OFFSET2:%[a-zA-Z0-9_]+8]])
diff --git a/clang/test/CodeGenSYCL/stream.cpp b/clang/test/CodeGenSYCL/stream.cpp
index 520f3c398b762..48ddfcc7cc6d9 100644
--- a/clang/test/CodeGenSYCL/stream.cpp
+++ b/clang/test/CodeGenSYCL/stream.cpp
@@ -5,7 +5,7 @@
 // CHECK: %[[ID_TYPE:"struct.*sycl::_V1::id"]]
 
 // CHECK: define dso_local spir_kernel void @{{.*}}StreamTester
-// CHECK-SAME: ptr addrspace(1) noundef align 1 [[ACC_DATA:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr addrspace(1) noundef [[ACC_DATA:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval(%[[ID_TYPE]]) align 4 [[ACC_ID:%[a-zA-Z0-9_]+]],