[SYCL] Use CUDA symbol for specialization constants

llvm/test/tools/sycl-post-link/spec-constants/composite-no-sret.ll

@@ -1,13 +1,13 @@
 ; RUN: sycl-post-link -spec-const=rt --ir-output-only < %s -S -o - \
 ; RUN: | FileCheck %s --implicit-check-not "call {{.*}} __sycl_getCompositeSpecConstantValue" --implicit-check-not "call {{.*}} __sycl_getComposite2020SpecConstantValue"
 
+; CHECK: declare i32 @_Z20__spirv_SpecConstantii(i32, i32)
+; CHECK: declare %struct._ZTS10TestStruct.TestStruct @_Z29__spirv_SpecConstantCompositeii_Rstruct._ZTS10TestStruct.TestStruct(i32, i32)
+
 ; CHECK: %[[#NS0:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID:]], i32
 ; CHECK: %[[#NS1:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 1]], i32 42)
 ; CHECK: %[[#NA0:]] = call %struct._ZTS10TestStruct.TestStruct @_Z29__spirv_SpecConstantCompositeii_Rstruct._ZTS10TestStruct.TestStruct(i32 %[[#NS0]], i32 %[[#NS1]])
 
-; CHECK: declare i32 @_Z20__spirv_SpecConstantii(i32, i32)
-; CHECK: declare %struct._ZTS10TestStruct.TestStruct @_Z29__spirv_SpecConstantCompositeii_Rstruct._ZTS10TestStruct.TestStruct(i32, i32)
-
 ; CHECK: !sycl.specialization-constants = !{![[#MD:]]}
 ; CHECK: ![[#MD]] = !{!"_ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL10SpecConst3EEE", i32 [[#ID]], i32 0, i32 4,
 

llvm/test/tools/sycl-post-link/spec-constants/cuda_spec_const.ll

@@ -0,0 +1,72 @@
+; RUN: sycl-post-link --ir-output-only --spec-const=rt %s -S -o - | FileCheck %s
+
+; This test is intended to check that CUDASpecConstantToSymbolPass is correctly
+; transforming the signature of the kernel and replacing accesses to the
+; implicit arg with the global variable.
+
+target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
+target triple = "nvptx64-nvidia-cuda"
+
+%"class.sycl::_V1::id" = type { %"class.sycl::_V1::detail::array" }
+%"class.sycl::_V1::detail::array" = type { [1 x i64] }
+%struct.spec_const_struct = type <{ i32, [4 x i8], i64, i8, [7 x i8] }>
+
+$_ZTS17spec_const_kernel = comdat any
+
+@__usid_str2 = private unnamed_addr constant [40 x i8] c"4a70cebf3f21eeb5____ZL15value_id_struct\00", align 1
+@_ZL15value_id_struct = internal addrspace(1) constant { { i32, i64, i8 } } { { i32, i64, i8 } { i32 1, i64 2, i8 3 } }, align 8
+
+; CHECK: sycl_specialization_constants_kernel__ZTS17spec_const_kernel
+
+; CHECK: void @_ZTS17spec_const_kernel
+; CHECK-NOT: %_arg__specialization_constants_buffer
+; Function Attrs: convergent noinline norecurse
+define weak_odr dso_local void @_ZTS17spec_const_kernel(i64 addrspace(1)* noundef align 8 %_arg_acc, %"class.sycl::_V1::id"* noundef byval(%"class.sycl::_V1::id") align 8 %_arg_acc3, i8 addrspace(1)* noundef align 1 %_arg__specialization_constants_buffer) local_unnamed_addr comdat {
+entry:
+  %0 = getelementptr inbounds %"class.sycl::_V1::id", %"class.sycl::_V1::id"* %_arg_acc3, i64 0, i32 0, i32 0, i64 0
+  %1 = load i64, i64* %0, align 8
+  %add.ptr.i = getelementptr inbounds i64, i64 addrspace(1)* %_arg_acc, i64 %1
+; CHECK-NOT: addrspacecast i8 addrspace(1)* %_arg__specialization_constants_buffer to i8*
+; CHECK: getelementptr i8, i8* addrspacecast (i8 addrspace(4)* getelementptr inbounds ([64 x i8], [64 x i8] addrspace(4)* @sycl_specialization_constants_kernel__ZTS17spec_const_kernel, i32 0, i32 0) to i8*), i32 0
+  %2 = addrspacecast i8 addrspace(1)* %_arg__specialization_constants_buffer to i8*
+  %gep = getelementptr i8, i8* %2, i32 0
+  %bc = bitcast i8* %gep to i64*
+  %load = load i64, i64* %bc, align 8
+; CHECK: getelementptr i8, i8* addrspacecast (i8 addrspace(4)* getelementptr inbounds ([64 x i8], [64 x i8] addrspace(4)* @sycl_specialization_constants_kernel__ZTS17spec_const_kernel, i32 0, i32 0) to i8*), i32 8
+  %gep1 = getelementptr i8, i8* %2, i32 8
+  %bc2 = bitcast i8* %gep1 to i8*
+  %load3 = load i8, i8* %bc2, align 1
+; CHECK: getelementptr i8, i8* addrspacecast (i8 addrspace(4)* getelementptr inbounds ([64 x i8], [64 x i8] addrspace(4)* @sycl_specialization_constants_kernel__ZTS17spec_const_kernel, i32 0, i32 0) to i8*), i32 9
+
+  %gep4 = getelementptr i8, i8* %2, i32 9
+  %bc5 = bitcast i8* %gep4 to %struct.spec_const_struct*
+  %load6 = load %struct.spec_const_struct, %struct.spec_const_struct* %bc5, align 1
+  %3 = extractvalue %struct.spec_const_struct %load6, 3
+  %conv.i = sext i8 %load3 to i64
+  %add.i = add nsw i64 %load, %conv.i
+  %conv4.i = sext i8 %3 to i64
+  %arrayidx.ascast.i.i = addrspacecast i64 addrspace(1)* %add.ptr.i to i64*
+  %4 = load i64, i64* %arrayidx.ascast.i.i, align 8
+  %add5.i = add i64 %add.i, %4
+  %add7.i = add i64 %add5.i, %conv4.i
+  store i64 %add7.i, i64* %arrayidx.ascast.i.i, align 8
+  ret void
+}
+
+!nvvm.annotations = !{!0}
+!opencl.spir.version = !{!1}
+!spirv.Source = !{!2}
+!llvm.module.flags = !{!3}
+!sycl.specialization-constants = !{!4}
+!sycl.specialization-constants-default-values = !{!5}
+!sycl.specialization-constants-kernel = !{!6}
+
+; CHECK-NOT: !{void (i64 addrspace(1)*, %"class.sycl::_V1::id"*, i8 addrspace(1)*)* @_ZTS17spec_const_kernel, !"kernel", i32 1}
+; CHECK: distinct !{void (i64 addrspace(1)*, %"class.sycl::_V1::id"*)* @_ZTS17spec_const_kernel, !"kernel", i32 1}
+!0 = !{void (i64 addrspace(1)*, %"class.sycl::_V1::id"*, i8 addrspace(1)*)* @_ZTS17spec_const_kernel, !"kernel", i32 1}
+!1 = !{i32 1, i32 2}
+!2 = !{i32 4, i32 100000}
+!3 = !{i32 1, !"wchar_size", i32 4}
+!4 = !{!"4a70cebf3f21eeb5____ZL15value_id_struct", i32 2, i32 0, i32 24}
+!5 = !{{ i32, i64, i8 } { i32 1, i64 2, i8 3 }}
+!6 = !{!"_ZTS17spec_const_kernel", !4}

llvm/test/tools/sycl-post-link/spec-constants/cuda_spec_const_no_accesses.ll

@@ -0,0 +1,44 @@
+; RUN: sycl-post-link --ir-output-only --spec-const=rt %s -S -o - | FileCheck %s
+
+; This test is intended to check that CUDASpecConstantToSymbolPass correctly
+; handles situations where _arg__specialization_constants_buffer is present,
+; however SpecConstants pass has not identified any uses of spec constants, and
+; hence both the implicit argument and the allocation should be present.
+
+target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
+target triple = "nvptx64-nvidia-cuda"
+
+%"class.sycl::_V1::id" = type { %"class.sycl::_V1::detail::array" }
+%"class.sycl::_V1::detail::array" = type { [1 x i64] }
+%struct.spec_const_struct = type <{ i32, [4 x i8], i64, i8, [7 x i8] }>
+
+$_ZTS17spec_const_kernel = comdat any
+
+@__usid_str2 = private unnamed_addr constant [40 x i8] c"4a70cebf3f21eeb5____ZL15value_id_struct\00", align 1
+@_ZL15value_id_struct = internal addrspace(1) constant { { i32, i64, i8 } } { { i32, i64, i8 } { i32 1, i64 2, i8 3 } }, align 8
+
+; CHECK: sycl_specialization_constants_kernel__ZTS17spec_const_kernel
+
+; CHECK: void @_ZTS17spec_const_kernel
+; CHECK: %_arg__specialization_constants_buffer
+; Function Attrs: convergent noinline norecurse
+define weak_odr dso_local void @_ZTS17spec_const_kernel(i64 addrspace(1)* noundef align 8 %_arg_acc, %"class.sycl::_V1::id"* noundef byval(%"class.sycl::_V1::id") align 8 %_arg_acc3, i8 addrspace(1)* noundef align 1 %_arg__specialization_constants_buffer) local_unnamed_addr comdat {
+entry:
+  ret void
+}
+
+!nvvm.annotations = !{!0}
+!opencl.spir.version = !{!1}
+!spirv.Source = !{!2}
+!llvm.module.flags = !{!3}
+!sycl.specialization-constants = !{!4}
+!sycl.specialization-constants-default-values = !{!5}
+
+; CHECK: !{void (i64 addrspace(1)*, %"class.sycl::_V1::id"*, i8 addrspace(1)*)* @_ZTS17spec_const_kernel, !"kernel", i32 1}
+!0 = !{void (i64 addrspace(1)*, %"class.sycl::_V1::id"*, i8 addrspace(1)*)* @_ZTS17spec_const_kernel, !"kernel", i32 1}
+!1 = !{i32 1, i32 2}
+!2 = !{i32 4, i32 100000}
+!3 = !{i32 1, !"wchar_size", i32 4}
+!4 = !{!"4a70cebf3f21eeb5____ZL15value_id_struct", i32 2, i32 0, i32 24}
+!5 = !{{ i32, i64, i8 } { i32 1, i64 2, i8 3 }}
+!6 = !{!"_ZTS17spec_const_kernel", !4}

llvm/test/tools/sycl-post-link/spec-constants/cuda_spec_const_wrong_triple.ll

@@ -0,0 +1,66 @@
+; RUN: sycl-post-link --ir-output-only --spec-const=rt %s -S -o - | FileCheck %s
+
+; This test is intended to check that CUDASpecConstantToSymbolPass does not
+; modify non-nvptx triples.
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+target triple = "spir64-unknown-unknown"
+
+%"class.sycl::_V1::id" = type { %"class.sycl::_V1::detail::array" }
+%"class.sycl::_V1::detail::array" = type { [1 x i64] }
+%struct.spec_const_struct = type <{ i32, [4 x i8], i64, i8, [7 x i8] }>
+
+$_ZTS17spec_const_kernel = comdat any
+
+@__usid_str2 = private unnamed_addr constant [40 x i8] c"4a70cebf3f21eeb5____ZL15value_id_struct\00", align 1
+@_ZL15value_id_struct = internal addrspace(1) constant { { i32, i64, i8 } } { { i32, i64, i8 } { i32 1, i64 2, i8 3 } }, align 8
+
+; CHECK-NOT: sycl_specialization_constants_kernel__ZTS17spec_const_kernel
+
+; CHECK: void @_ZTS17spec_const_kernel
+; CHECK: %_arg__specialization_constants_buffer
+; Function Attrs: convergent noinline norecurse
+define weak_odr dso_local void @_ZTS17spec_const_kernel(i64 addrspace(1)* noundef align 8 %_arg_acc, %"class.sycl::_V1::id"* noundef byval(%"class.sycl::_V1::id") align 8 %_arg_acc3, i8 addrspace(1)* noundef align 1 %_arg__specialization_constants_buffer) local_unnamed_addr comdat {
+entry:
+  %0 = getelementptr inbounds %"class.sycl::_V1::id", %"class.sycl::_V1::id"* %_arg_acc3, i64 0, i32 0, i32 0, i64 0
+  %1 = load i64, i64* %0, align 8
+  %add.ptr.i = getelementptr inbounds i64, i64 addrspace(1)* %_arg_acc, i64 %1
+; CHECK: addrspacecast i8 addrspace(1)* %_arg__specialization_constants_buffer to i8*
+  %2 = addrspacecast i8 addrspace(1)* %_arg__specialization_constants_buffer to i8*
+  %gep = getelementptr i8, i8* %2, i32 0
+  %bc = bitcast i8* %gep to i64*
+  %load = load i64, i64* %bc, align 8
+  %gep1 = getelementptr i8, i8* %2, i32 8
+  %bc2 = bitcast i8* %gep1 to i8*
+  %load3 = load i8, i8* %bc2, align 1
+  %gep4 = getelementptr i8, i8* %2, i32 9
+  %bc5 = bitcast i8* %gep4 to %struct.spec_const_struct*
+  %load6 = load %struct.spec_const_struct, %struct.spec_const_struct* %bc5, align 1
+  %3 = extractvalue %struct.spec_const_struct %load6, 3
+  %conv.i = sext i8 %load3 to i64
+  %add.i = add nsw i64 %load, %conv.i
+  %conv4.i = sext i8 %3 to i64
+  %arrayidx.ascast.i.i = addrspacecast i64 addrspace(1)* %add.ptr.i to i64*
+  %4 = load i64, i64* %arrayidx.ascast.i.i, align 8
+  %add5.i = add i64 %add.i, %4
+  %add7.i = add i64 %add5.i, %conv4.i
+  store i64 %add7.i, i64* %arrayidx.ascast.i.i, align 8
+  ret void
+}
+
+!nvvm.annotations = !{!0}
+!opencl.spir.version = !{!1}
+!spirv.Source = !{!2}
+!llvm.module.flags = !{!3}
+!sycl.specialization-constants = !{!4}
+!sycl.specialization-constants-default-values = !{!5}
+!sycl.specialization-constants-kernel = !{!6}
+
+; CHECK: !{void (i64 addrspace(1)*, %"class.sycl::_V1::id"*, i8 addrspace(1)*)* @_ZTS17spec_const_kernel, !"kernel", i32 1}
+!0 = !{void (i64 addrspace(1)*, %"class.sycl::_V1::id"*, i8 addrspace(1)*)* @_ZTS17spec_const_kernel, !"kernel", i32 1}
+!1 = !{i32 1, i32 2}
+!2 = !{i32 4, i32 100000}
+!3 = !{i32 1, !"wchar_size", i32 4}
+!4 = !{!"4a70cebf3f21eeb5____ZL15value_id_struct", i32 2, i32 0, i32 24}
+!5 = !{{ i32, i64, i8 } { i32 1, i64 2, i8 3 }}
+!6 = !{!"_ZTS17spec_const_kernel", !4}

llvm/tools/sycl-post-link/CMakeLists.txt

@@ -25,6 +25,7 @@ add_llvm_tool(sycl-post-link
   sycl-post-link.cpp
   ModuleSplitter.cpp
   SpecConstants.cpp
+  CUDASpecConstantToSymbol.cpp
   SYCLDeviceLibReqMask.cpp
   SYCLKernelParamOptInfo.cpp
   SYCLDeviceRequirements.cpp