22// UNSUPPORTED: system-windows
33// RUN: dpct --no-dpcpp-extensions=device_info --format-range=none --usm-level=none -out-root %T/memory_management %s --cuda-include-path="%cuda-path/include" -output-file=memory_management_outputfile.txt -- -x cuda --cuda-host-only
44// RUN: FileCheck --match-full-lines --input-file %T/memory_management/memory_management.dp.cpp %s
5+ // RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/memory_management/memory_management.dp.cpp -o %T/memory_management/memory_management.dp.o %}
56
7+ #ifndef BUILD_TEST
8+
9+ #include < cuda.h>
610#include < cuda_runtime.h>
7- #include < cuda.h>
811
912__constant__ float constData[123 * 4 ];
1013
11- // CHECK: template<typename T>
12- // CHECK-NEXT: void test(){
13- // CHECK-NEXT: int i = 0;
14- // CHECK-NEXT: T** ptr;
15- // CHECK-NEXT: T* array[10];
16- // CHECK-NEXT: ptr[i] = (T *)dpct::dpct_malloc(10 * sizeof(T));
17- // CHECK-NEXT: ptr[i] = (T *)dpct::dpct_malloc(10 * sizeof(T));
18- // CHECK-NEXT: array[i] = (T *)dpct::dpct_malloc(10 * sizeof(T));
19- // CHECK-NEXT: }
20- template <typename T>
21- void test (){
14+ // CHECK: template <typename T>
15+ // CHECK-NEXT: void test() {
16+ // CHECK-NEXT: int i = 0;
17+ // CHECK-NEXT: T ** ptr;
18+ // CHECK-NEXT: T * array[10];
19+ // CHECK-NEXT: ptr[i] = (T *)dpct::dpct_malloc(10 * sizeof(T));
20+ // CHECK-NEXT: ptr[i] = (T *)dpct::dpct_malloc(10 * sizeof(T));
21+ // CHECK-NEXT: array[i] = (T *)dpct::dpct_malloc(10 * sizeof(T));
22+ // CHECK-NEXT: }
23+ template <typename T>
24+ void test () {
2225 int i = 0 ;
23- T** ptr;
24- T* array[10 ];
26+ T ** ptr;
27+ T * array[10 ];
2528 cudaMalloc (&ptr[i], 10 * sizeof (T));
2629 cudaMalloc (&(ptr[i]), 10 * sizeof (T));
2730 cudaMalloc (&array[i], 10 * sizeof (T));
@@ -195,8 +198,8 @@ void checkError(cudaError_t err) {
195198
196199void cuCheckError (CUresult err) {
197200}
198- // CHECK: #define PITCH(a,b,c, d) a = (float *)dpct::dpct_malloc(b, c, d);
199- #define PITCH (a,b,c, d ) cudaMallocPitch(a, b, c, d);
201+ // CHECK: #define PITCH(a, b, c, d) a = (float *)dpct::dpct_malloc(b, c, d);
202+ #define PITCH (a, b, c, d ) cudaMallocPitch(a, b, c, d);
200203
201204void testCommas () {
202205 size_t size = 1234567 * sizeof (float );
@@ -426,7 +429,6 @@ void testCommas() {
426429 // CHECK: checkError(DPCT_CHECK_ERROR(dpct::dpct_memcpy(d_B, constData.get_ptr(), size, dpct::device_to_device)));
427430 checkError (cudaMemcpyFromSymbol (d_B, constData, size, 0 , cudaMemcpyDeviceToDevice));
428431
429-
430432 // CHECK: dpct::dpct_memcpy(d_B, (char *)(constData.get_ptr()) + 1, size, dpct::device_to_device);
431433 cudaMemcpyFromSymbol (d_B, constData, size, 1 , cudaMemcpyDeviceToDevice);
432434 // CHECK: dpct::dpct_memcpy(d_B, (char *)(constData.get_ptr()) + 1, size, dpct::device_to_device);
@@ -614,7 +616,6 @@ void testCommas_in_global_memory() {
614616 // CHECK: result2 = dpct::get_current_device().get_device_info().get_global_mem_size();
615617 cudaMemGetInfo (&result1, &result2);
616618
617-
618619 // CHECK: /*
619620 // CHECK: DPCT1072:{{[0-9]+}}: SYCL currently does not support getting the available memory on the current device. You may need to adjust the code.
620621 // CHECK: */
@@ -627,7 +628,7 @@ void testCommas_in_global_memory() {
627628 // CHECK: checkError(DPCT_CHECK_ERROR(result2 = dpct::get_current_device().get_device_info().get_global_mem_size()));
628629 checkError (cudaMemGetInfo (&result1, &result2));
629630
630- CUdeviceptr devicePtr;
631+ CUdeviceptr devicePtr;
631632 // CHECK: devicePtr = (dpct::device_ptr)dpct::dpct_malloc(size, size, size);
632633 cuMemAllocPitch ((CUdeviceptr *)&devicePtr, &size, size, size, size);
633634
@@ -636,15 +637,15 @@ void testCommas_in_global_memory() {
636637 // CHECK: cuCheckError(DPCT_CHECK_ERROR(devicePtr = (dpct::device_ptr)dpct::dpct_malloc(size, size, size)));
637638 cuCheckError (cuMemAllocPitch ((CUdeviceptr *)&devicePtr, &size, size, size, size));
638639
639- int * a;
640+ int * a;
640641 cudaStream_t stream;
641642 int deviceID = 0 ;
642- CUdevice cudevice =0 ;
643+ CUdevice cudevice = 0 ;
643644 CUdeviceptr devPtr;
644645 // CHECK:/*
645646 // CHECK-NEXT:DPCT1007:{{[0-9]+}}: Migration of cudaMemPrefetchAsync is not supported.
646647 // CHECK-NEXT:*/
647- cudaMemPrefetchAsync (a, 100 , deviceID, stream);
648+ cudaMemPrefetchAsync (a, 100 , deviceID, stream);
648649
649650 // CHECK:/*
650651 // CHECK-NEXT:DPCT1007:{{[0-9]+}}: Migration of cuMemPrefetchAsync is not supported.
@@ -655,34 +656,33 @@ void testCommas_in_global_memory() {
655656 free (h_A);
656657}
657658
658- #define MY_CHECKER (CALL ) \
659- if ((CALL) != cudaSuccess) { \
660- exit (-1 ); \
661- }
659+ #define MY_CHECKER (CALL ) \
660+ if ((CALL) != cudaSuccess) { \
661+ exit (-1 ); \
662+ }
662663
663664#define MY_ERROR_CHECKER (CALL ) my_error_checker((CALL), #CALL)
664665template <typename T>
665666void my_error_checker (T ReturnValue, char const *const FuncName) {}
666667
667-
668- template <typename T>
669- void uninstantiated_template_call (const T * d_data, size_t width, size_t height) {
668+ template <typename T>
669+ void uninstantiated_template_call (const T *d_data, size_t width, size_t height) {
670670 size_t datasize = width * height;
671- T * data = new T[datasize];
671+ T *data = new T[datasize];
672672 cudaMemcpy3DParms parms;
673673 // CHECK: assert_cuda(DPCT_CHECK_ERROR(dpct::dpct_memcpy(data, d_data, datasize * sizeof(T), dpct::device_to_host)));
674674 assert_cuda (cudaMemcpy (data, d_data, datasize * sizeof (T), cudaMemcpyDeviceToHost));
675675
676676 // CHECK: dpct::dpct_memcpy(data, d_data, datasize * sizeof(T), dpct::device_to_host);
677677 cudaMemcpy (data, d_data, datasize * sizeof (T), cudaMemcpyDeviceToHost);
678678
679- #define DATAMACRO data+ 32 * 32
679+ #define DATAMACRO data + 32 * 32
680680
681681 // CHECK: dpct::dpct_memcpy(DATAMACRO, d_data, datasize * sizeof(T), dpct::device_to_host);
682682 cudaMemcpy (DATAMACRO, d_data, datasize * sizeof (T), cudaMemcpyDeviceToHost);
683683
684- // CHECK: dpct::dpct_memcpy(32*32+ DATAMACRO, d_data, datasize * sizeof(T), dpct::device_to_host);
685- cudaMemcpy (32 * 32 + DATAMACRO, d_data, datasize * sizeof (T), cudaMemcpyDeviceToHost);
684+ // CHECK: dpct::dpct_memcpy(32 * 32 + DATAMACRO, d_data, datasize * sizeof(T), dpct::device_to_host);
685+ cudaMemcpy (32 * 32 + DATAMACRO, d_data, datasize * sizeof (T), cudaMemcpyDeviceToHost);
686686
687687 // CHECK: checkError(DPCT_CHECK_ERROR(dpct::dpct_memcpy(data, d_data, datasize * sizeof(T), dpct::device_to_host)));
688688 checkError (cudaMemcpy (data, d_data, datasize * sizeof (T), cudaMemcpyDeviceToHost));
@@ -696,9 +696,9 @@ void uninstantiated_template_call(const T * d_data, size_t width, size_t height)
696696 // CHECK: MY_ERROR_CHECKER(DPCT_CHECK_ERROR(dpct::dpct_memcpy(data, d_data, datasize * sizeof(T), dpct::device_to_host)));
697697 MY_ERROR_CHECKER (cudaMemcpy (data, d_data, datasize * sizeof (T), cudaMemcpyDeviceToHost));
698698
699- // CHECK: #define CUDAMEMCPY dpct::dpct_memcpy
700- // CHECK-NEXT: CUDAMEMCPY(data, d_data, datasize * sizeof(T), dpct::device_to_host);
701- #define CUDAMEMCPY cudaMemcpy
699+ // CHECK: #define CUDAMEMCPY dpct::dpct_memcpy
700+ // CHECK-NEXT: CUDAMEMCPY(data, d_data, datasize * sizeof(T), dpct::device_to_host);
701+ #define CUDAMEMCPY cudaMemcpy
702702 CUDAMEMCPY (data, d_data, datasize * sizeof (T), cudaMemcpyDeviceToHost);
703703
704704 // CHECK: assert_cuda(DPCT_CHECK_ERROR(dpct::dpct_memcpy(data, datasize, d_data, datasize, datasize, datasize, dpct::device_to_host)));
@@ -710,18 +710,18 @@ void uninstantiated_template_call(const T * d_data, size_t width, size_t height)
710710 // CHECK: dpct::dpct_memcpy(DATAMACRO, datasize, d_data, datasize, datasize, datasize, dpct::device_to_host);
711711 cudaMemcpy2D (DATAMACRO, datasize, d_data, datasize, datasize, datasize, cudaMemcpyDeviceToHost);
712712
713- // CHECK: dpct::dpct_memcpy(32*32+ DATAMACRO, datasize, d_data, datasize, datasize, datasize, dpct::device_to_host);
714- cudaMemcpy2D (32 * 32 + DATAMACRO, datasize, d_data, datasize, datasize, datasize, cudaMemcpyDeviceToHost);
713+ // CHECK: dpct::dpct_memcpy(32 * 32 + DATAMACRO, datasize, d_data, datasize, datasize, datasize, dpct::device_to_host);
714+ cudaMemcpy2D (32 * 32 + DATAMACRO, datasize, d_data, datasize, datasize, datasize, cudaMemcpyDeviceToHost);
715715
716716 // CHECK: MY_CHECKER(DPCT_CHECK_ERROR(dpct::dpct_memcpy(data, datasize, d_data, datasize, datasize, datasize, dpct::device_to_host)));
717717 MY_CHECKER (cudaMemcpy2D (data, datasize, d_data, datasize, datasize, datasize, cudaMemcpyDeviceToHost));
718718
719719 // CHECK: MY_ERROR_CHECKER(DPCT_CHECK_ERROR(dpct::dpct_memcpy(data, datasize, d_data, datasize, datasize, datasize, dpct::device_to_host)));
720720 MY_ERROR_CHECKER (cudaMemcpy2D (data, datasize, d_data, datasize, datasize, datasize, cudaMemcpyDeviceToHost));
721721
722- // CHECK: #define CUDAMEMCPY2D dpct::dpct_memcpy
723- // CHECK-NEXT: CUDAMEMCPY2D(data, datasize, d_data, datasize, datasize, datasize, dpct::device_to_host);
724- #define CUDAMEMCPY2D cudaMemcpy2D
722+ // CHECK: #define CUDAMEMCPY2D dpct::dpct_memcpy
723+ // CHECK-NEXT: CUDAMEMCPY2D(data, datasize, d_data, datasize, datasize, datasize, dpct::device_to_host);
724+ #define CUDAMEMCPY2D cudaMemcpy2D
725725 CUDAMEMCPY2D (data, datasize, d_data, datasize, datasize, datasize, cudaMemcpyDeviceToHost);
726726
727727 // CHECK: MY_CHECKER(DPCT_CHECK_ERROR(dpct::dpct_memcpy(parms)));
@@ -737,18 +737,18 @@ void uninstantiated_template_call(const T * d_data, size_t width, size_t height)
737737void test_segmentation_fault () {
738738 float *buffer;
739739 /*
740- * Original code in getSizeString():
741- * "SizeExpr->getBeginLoc()" cannot get the real SourceLocation of "N*sizeof(float)",
742- * and results in boundary violation in "dpctGlobalInfo::getSourceManager().getCharacterData(SizeBegin)"
743- * and fails with segmentation fault.
744- */
745- cudaMalloc (&buffer, N* sizeof (float ));
740+ * Original code in getSizeString():
741+ * "SizeExpr->getBeginLoc()" cannot get the real SourceLocation of "N*sizeof(float)",
742+ * and results in boundary violation in "dpctGlobalInfo::getSourceManager().getCharacterData(SizeBegin)"
743+ * and fails with segmentation fault.
744+ */
745+ cudaMalloc (&buffer, N * sizeof (float ));
746746}
747747
748748// CHECK: static dpct::global_memory<uint32_t, 1> d_error(1);
749749static __device__ uint32_t d_error[1 ];
750750
751- void test_foo (){
751+ void test_foo () {
752752 // CHECK: dpct::dpct_memset(d_error.get_ptr(), 0, sizeof(uint32_t));
753753 cudaMemset (d_error, 0 , sizeof (uint32_t ));
754754}
@@ -766,18 +766,18 @@ void foobar() {
766766 // CHECK: flags = 0;
767767 cudaArrayGetInfo (&desc, &extent, &flags, array);
768768
769- // CHECK: checkError(DPCT_CHECK_ERROR([&](){
770- // CHECK-NEXT: desc = array->get_channel();
771- // CHECK-NEXT: extent = array->get_range();
772- // CHECK-NEXT: flags = 0;
773- // CHECK-NEXT: }()));
769+ // CHECK: checkError(DPCT_CHECK_ERROR([&](){
770+ // CHECK-NEXT: desc = array->get_channel();
771+ // CHECK-NEXT: extent = array->get_range();
772+ // CHECK-NEXT: flags = 0;
773+ // CHECK-NEXT: }()));
774774 checkError (cudaArrayGetInfo (&desc, &extent, &flags, array));
775775
776- // CHECK: errorCode = DPCT_CHECK_ERROR([&](){
777- // CHECK-NEXT: desc = array->get_channel();
778- // CHECK-NEXT: extent = array->get_range();
779- // CHECK-NEXT: flags = 0;
780- // CHECK-NEXT: }());
776+ // CHECK: errorCode = DPCT_CHECK_ERROR([&](){
777+ // CHECK-NEXT: desc = array->get_channel();
778+ // CHECK-NEXT: extent = array->get_range();
779+ // CHECK-NEXT: flags = 0;
780+ // CHECK-NEXT: }());
781781 errorCode = cudaArrayGetInfo (&desc, &extent, &flags, array);
782782
783783 int host;
@@ -803,7 +803,6 @@ void foobar() {
803803 */
804804 CUmemAccessDesc c;
805805
806-
807806 int *devPtr;
808807
809808 CUdeviceptr devicePtr;
@@ -812,7 +811,7 @@ void foobar() {
812811
813812 CUdeviceptr cuDevPtr;
814813
815- CUdevice cudevice =0 ;
814+ CUdevice cudevice = 0 ;
816815
817816 CUmem_advise advise = CU_MEM_ADVISE_UNSET_PREFERRED_LOCATION;
818817
@@ -867,17 +866,18 @@ void foobar() {
867866 errorCode = cudaMemAdvise (devPtr, count, cudaMemAdviseSetReadMostly, device);
868867}
869868
870- // CHECK: void copy_dir_1 (dpct::memcpy_direction kind) {}
871- // CHECK-NEXT: void copy_dir_2 (dpct::memcpy_direction kind) {}
872- // CHECK-NEXT: void copy_dir_3 (dpct::memcpy_direction kind) {}
873- void copy_dir_1 (cudaMemcpyKind kind) {}
874- void copy_dir_2 (enum cudaMemcpyKind kind) {}
875- void copy_dir_3 (enum cudaMemcpyKind kind) {}
876-
877- // CHECK: void copy_dir_1 (int kind) {}
878- // CHECK-NEXT: void copy_dir_2 (int kind) {}
879- // CHECK-NEXT: void copy_dir_3 (int kind) {}
880- void copy_dir_1 (cudaComputeMode kind) {}
881- void copy_dir_2 (enum cudaComputeMode kind) {}
882- void copy_dir_3 (enum cudaComputeMode kind) {}
883-
869+ // CHECK: void copy_dir_1(dpct::memcpy_direction kind) {}
870+ // CHECK-NEXT: void copy_dir_2(dpct::memcpy_direction kind) {}
871+ // CHECK-NEXT: void copy_dir_3(dpct::memcpy_direction kind) {}
872+ void copy_dir_1 (cudaMemcpyKind kind) {}
873+ void copy_dir_2 (enum cudaMemcpyKind kind) {}
874+ void copy_dir_3 (enum cudaMemcpyKind kind) {}
875+
876+ // CHECK: void copy_dir_1(int kind) {}
877+ // CHECK-NEXT: void copy_dir_2(int kind) {}
878+ // CHECK-NEXT: void copy_dir_3(int kind) {}
879+ void copy_dir_1 (cudaComputeMode kind) {}
880+ void copy_dir_2 (enum cudaComputeMode kind) {}
881+ void copy_dir_3 (enum cudaComputeMode kind) {}
882+
883+ #endif
0 commit comments