From 05965deab59b93f6142e8f6ff4d7d66a90394077 Mon Sep 17 00:00:00 2001 From: Teja Alaghari Date: Tue, 9 Jul 2024 19:10:08 +0800 Subject: [PATCH 01/10] Enabled migration of cuPointerGetAttributes --- .../DPCT/Driver/cuPointerGetAttributes.cu | 9 ++ clang/lib/DPCT/APINamesMemory.inc | 7 ++ clang/lib/DPCT/ASTTraversal.cpp | 33 ++--- clang/lib/DPCT/MapNames.cpp | 46 +++++++ clang/runtime/dpct-rt/include/dpct/memory.hpp | 53 ++++++++ clang/test/dpct/cuPointerGetAttributes.cu | 119 ++++++++++++++++++ .../dpct/query_api_mapping/Driver/test.cu | 11 ++ clang/test/dpct/query_api_mapping/test_all.cu | 1 + ...untime_and_Driver_API_migration_status.csv | 2 +- 9 files changed, 265 insertions(+), 16 deletions(-) create mode 100644 clang/examples/DPCT/Driver/cuPointerGetAttributes.cu create mode 100644 clang/test/dpct/cuPointerGetAttributes.cu diff --git a/clang/examples/DPCT/Driver/cuPointerGetAttributes.cu b/clang/examples/DPCT/Driver/cuPointerGetAttributes.cu new file mode 100644 index 000000000000..921fc777d1d0 --- /dev/null +++ b/clang/examples/DPCT/Driver/cuPointerGetAttributes.cu @@ -0,0 +1,9 @@ +void test(unsigned int numAttributes, CUpointer_attribute *attributes, + void **data, CUdeviceptr ptr) { + // Start + cuPointerGetAttribute(numAttributes /*unsigned int*/, + attributes /*CUpointer_attribute **/, data /*void ***/, + ptr /*CUdeviceptr*/); + // End +} + diff --git a/clang/lib/DPCT/APINamesMemory.inc b/clang/lib/DPCT/APINamesMemory.inc index 8e353282a72e..1350d3a0ccde 100644 --- a/clang/lib/DPCT/APINamesMemory.inc +++ b/clang/lib/DPCT/APINamesMemory.inc @@ -398,6 +398,13 @@ ASSIGNABLE_FACTORY( MEMBER_CALL_FACTORY_ENTRY( "cudaPointerGetAttributes", DEREF(0), false, "init", ARG(1)))) +ASSIGNABLE_FACTORY( + FEATURE_REQUEST_FACTORY( + HelperFeatureEnum::device_ext, + CALL_FACTORY_ENTRY( + "cuPointerGetAttributes", CALL(MapNames::getDpctNamespace() + "pointer_attributes.get", + ARG(0), ARG(1), ARG(2), ARG(3))))) + ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( HelperFeatureEnum::device_ext, CONDITIONAL_FACTORY_ENTRY( diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index 46c79fd908d9..3e73f2b51919 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -1727,9 +1727,10 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) { "__nv_bfloat16", "cooperative_groups::__v1::thread_group", "cooperative_groups::__v1::thread_block", "libraryPropertyType_t", "libraryPropertyType", "cudaDataType_t", "cudaDataType", - "cublasComputeType_t", "cublasAtomicsMode_t", "CUmem_advise_enum", - "CUmem_advise", "thrust::tuple_element", "thrust::tuple_size", - "cublasMath_t", "cudaPointerAttributes", "thrust::zip_iterator", + "cublasComputeType_t", "cublasAtomicsMode_t", "cublasMath_t", + "CUmem_advise_enum", "CUmem_advise", "thrust::tuple_element", + "thrust::tuple_size", "thrust::zip_iterator", + "cudaPointerAttributes", "CUpointer_attribute", "cusolverEigRange_t", "cudaUUID_t", "cusolverDnFunction_t", "cusolverAlgMode_t", "cusparseIndexType_t", "cusparseFormat_t", "cusparseDnMatDescr_t", "cusparseOrder_t", "cusparseDnVecDescr_t", @@ -3210,14 +3211,15 @@ REGISTER_RULE(DeviceInfoVarRule, PassKind::PK_Migration) void EnumConstantRule::registerMatcher(MatchFinder &MF) { MF.addMatcher( declRefExpr( - to(enumConstantDecl(anyOf( - hasType(enumDecl(hasAnyName( - "cudaComputeMode", "cudaMemcpyKind", "cudaMemoryAdvise", - "cudaStreamCaptureStatus", "cudaDeviceAttr", - "libraryPropertyType_t", "cudaDataType_t", - "cublasComputeType_t", "CUmem_advise_enum", "cufftType_t", - "cufftType", "cudaMemoryType", "CUctx_flags_enum"))), - matchesName("CUDNN_.*"), matchesName("CUSOLVER_.*"))))) + to(enumConstantDecl( + anyOf(hasType(enumDecl(hasAnyName( + "cudaComputeMode", "cudaMemcpyKind", "cudaMemoryAdvise", + "cudaStreamCaptureStatus", "cudaDeviceAttr", + "libraryPropertyType_t", "cudaDataType_t", + "cublasComputeType_t", "CUmem_advise_enum", + "cufftType_t", "cufftType", "cudaMemoryType", + "CUctx_flags_enum", "CUpointer_attribute_enum"))), + matchesName("CUDNN_.*"), matchesName("CUSOLVER_.*"))))) .bind("EnumConstant"), this); } @@ -6109,10 +6111,11 @@ void FunctionCallRule::registerMatcher(MatchFinder &MF) { "cuDeviceCanAccessPeer", "cudaFuncSetAttribute", "cudaRuntimeGetVersion", "clock64", "__nanosleep", "cudaFuncSetSharedMemConfig", "cuFuncSetCacheConfig", - "cudaPointerGetAttributes", "cuCtxSetCacheConfig", "cuCtxSetLimit", - "cudaCtxResetPersistingL2Cache", "cuCtxResetPersistingL2Cache", - "cudaStreamSetAttribute", "cudaStreamGetAttribute", "cudaProfilerStart", - "cudaProfilerStop", "__trap", "cuCtxEnablePeerAccess"); + "cudaPointerGetAttributes", "cuPointerGetAttributes", + "cuCtxSetCacheConfig", "cuCtxSetLimit", "cudaCtxResetPersistingL2Cache", + "cuCtxResetPersistingL2Cache", "cudaStreamSetAttribute", + "cudaStreamGetAttribute", "cudaProfilerStart", "cudaProfilerStop", + "__trap", "cuCtxEnablePeerAccess"); }; MF.addMatcher( diff --git a/clang/lib/DPCT/MapNames.cpp b/clang/lib/DPCT/MapNames.cpp index 53185421f7f8..fd8879215e09 100644 --- a/clang/lib/DPCT/MapNames.cpp +++ b/clang/lib/DPCT/MapNames.cpp @@ -108,6 +108,8 @@ void MapNames::setExplicitNamespaceMap() { {"CUfunction", std::make_shared(getDpctNamespace() + "kernel_function", HelperFeatureEnum::device_ext)}, + {"CUpointer_attribute", + std::make_shared(getDpctNamespace() + "attribute_type")}, {"cudaPointerAttributes", std::make_shared(getDpctNamespace() + "pointer_attributes", HelperFeatureEnum::device_ext)}, @@ -1175,6 +1177,50 @@ void MapNames::setExplicitNamespaceMap() { std::make_shared("get_device_info().get_local_mem_size", HelperFeatureEnum::device_ext)}, + // enum CUpointer_attribute + {"CU_POINTER_ATTRIBUTE_CONTEXT", + std::make_shared(getDpctNamespace() + + "attribute_type::context")}, + {"CU_POINTER_ATTRIBUTE_MEMORY_TYPE", + std::make_shared(getDpctNamespace() + + "attribute_type::memory_type")}, + {"CU_POINTER_ATTRIBUTE_DEVICE_POINTER", + std::make_shared(getDpctNamespace() + + "attribute_type::device_pointer")}, + {"CU_POINTER_ATTRIBUTE_HOST_POINTER", + std::make_shared(getDpctNamespace() + + "attribute_type::host_pointer")}, + {"CU_POINTER_ATTRIBUTE_P2P_TOKENS", + std::make_shared(getDpctNamespace() + + "attribute_type::p2p_tokens")}, + {"CU_POINTER_ATTRIBUTE_SYNC_MEMOPS", + std::make_shared(getDpctNamespace() + + "attribute_type::sync_memps")}, + {"CU_POINTER_ATTRIBUTE_BUFFER_ID", + std::make_shared(getDpctNamespace() + + "attribute_type::buffer_id")}, + {"CU_POINTER_ATTRIBUTE_IS_MANAGED", + std::make_shared(getDpctNamespace() + + "attribute_type::is_managed")}, + {"CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL", + std::make_shared(getDpctNamespace() + + "attribute_type::device_id")}, + {"CU_POINTER_ATTRIBUTE_IS_LEGACY_CUDA_IPC_CAPABLE", + std::make_shared(getDpctNamespace() + + "attribute_type::is_legacy_ipc_capable")}, + {"CU_POINTER_ATTRIBUTE_RANGE_START_ADDR", + std::make_shared(getDpctNamespace() + + "attribute_type::range_start_addr")}, + {"CU_POINTER_ATTRIBUTE_RANGE_SIZE", + std::make_shared(getDpctNamespace() + + "attribute_type::range_size")}, + {"CU_POINTER_ATTRIBUTE_MAPPED", + std::make_shared(getDpctNamespace() + + "attribute_type::mapped")}, + {"CU_POINTER_ATTRIBUTE_ALLOWED_HANDLE_TYPES", + std::make_shared(getDpctNamespace() + + "attribute_type::allowed_handle_types")}, + // enum CUlimit {"CU_LIMIT_PRINTF_FIFO_SIZE", std::make_shared("INT_MAX")}, diff --git a/clang/runtime/dpct-rt/include/dpct/memory.hpp b/clang/runtime/dpct-rt/include/dpct/memory.hpp index f60e16ac9b23..94a33e77b717 100644 --- a/clang/runtime/dpct-rt/include/dpct/memory.hpp +++ b/clang/runtime/dpct-rt/include/dpct/memory.hpp @@ -1653,6 +1653,23 @@ template using usm_device_allocator = detail::deprecated::usm_allocator; } // namespace deprecated +enum class attribute_type { + context, + memory_type, + device_pointer, + host_pointer, + p2p_tokens, + sync_memps, + buffer_id, + is_managed, + device_id, + is_legacy_ipc_capable, + range_start_addr, + range_size, + mapped, + allowed_handle_types +}; + class pointer_attributes { public: void init(const void *ptr, @@ -1676,6 +1693,41 @@ class pointer_attributes { #endif } + static void get(unsigned int numAttributes, + attribute_type *attributes, void **data, + device_ptr ptr) { + pointer_attributes sycl_attributes; + + sycl_attributes.init(ptr); + + for (int i = 0; i < numAttributes; i++) { + switch (attributes[i]) { + case attribute_type::memory_type: + *static_cast(data[i]) = + static_cast(sycl_attributes.get_memory_type()); + break; + case attribute_type::device_pointer: + *(reinterpret_cast(data[i])) = + const_cast(sycl_attributes.get_device_pointer()); + break; + case attribute_type::host_pointer: + *(reinterpret_cast(data[i])) = + const_cast(sycl_attributes.get_host_pointer()); + break; + case attribute_type::is_managed: + *static_cast(data[i]) = + sycl_attributes.is_memory_shared(); + break; + case attribute_type::device_id: + *static_cast(data[i]) = sycl_attributes.get_device_id(); + break; + default: + data[i] = nullptr; + break; + } + } + } + sycl::usm::alloc get_memory_type() { return memory_type; } @@ -1702,5 +1754,6 @@ class pointer_attributes { const void *host_pointer = nullptr; unsigned int device_id = -1; }; + } // namespace dpct #endif // __DPCT_MEMORY_HPP__ diff --git a/clang/test/dpct/cuPointerGetAttributes.cu b/clang/test/dpct/cuPointerGetAttributes.cu new file mode 100644 index 000000000000..309ef958ade0 --- /dev/null +++ b/clang/test/dpct/cuPointerGetAttributes.cu @@ -0,0 +1,119 @@ +// UNSUPPORTED: v11.0, v11.1, v11.2, v11.3, v11.4, v11.5, v11.6, v11.7, v11.8, v12.0, v12.1, v12.2, v12.3 , v12.4 +// UNSUPPORTED: cuda-11.0, cuda-11.1, cuda-11.2, cuda-11.3, cuda-11.4, cuda-11.5, cuda-11.6, cuda-11.7, cuda-11.8, cuda-12.0, cuda-12.1, cuda-12.2, cuda-12.3, cuda-12.4 +// RUN: dpct --format-range=none -out-root %T/Out/cudaPointerAttributes %s --cuda-include-path="%cuda-path/include" +// RUN: FileCheck %s --match-full-lines --input-file %T/Out/cudaPointerAttributes/cudaPointerAttributes.dp.cpp +// RUN: %if build_lit %{icpx -c -fsycl %T/Out/cudaPointerAttributes/cudaPointerAttributes.dp.cpp -o %T/Out/cudaPointerAttributes/cudaPointerAttributes.dp.o %} +#include +#include +#include +#include +int main() { + int N = 2048; + size_t size = N * sizeof(float); + + float *h_A = (float *)malloc(size); + + float *d_A; + cudaMalloc((void **)&d_A, size); + + cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); + cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice); + + unsigned int numAttributes = 5; + + // CHECK: dpct::attribute_type attributes[] = { + // CHECK: dpct::attribute_type::memory_type, dpct::attribute_type::device_pointer, + // CHECK: dpct::attribute_type::host_pointer, dpct::attribute_type::is_managed, + // CHECK: dpct::attribute_type::device_id}; + CUpointer_attribute attributes[] = { + CU_POINTER_ATTRIBUTE_MEMORY_TYPE, + CU_POINTER_ATTRIBUTE_DEVICE_POINTER, + CU_POINTER_ATTRIBUTE_HOST_POINTER, + CU_POINTER_ATTRIBUTE_IS_MANAGED, + CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL + }; + + int memType; + // CHECK: dpct::device_ptr devPtr; + CUdeviceptr devPtr; + void* hostPtr; + unsigned int isManaged; + int deviceID; + + void* attributeValues[] = { + &memType, + &devPtr, + &hostPtr, + &isManaged, + &deviceID + }; + + // CHECK: dpct::pointer_attributes.get( + cuPointerGetAttributes( + numAttributes, + attributes, + attributeValues, + h_A + ); + + std::cout << "====== Host Attributes =======" << std::endl; + std::cout << deviceID << std::endl; + std::cout << memType << std::endl; + std::cout << hostPtr << std::endl; + std::cout << devPtr << std::endl; + std::cout << isManaged << std::endl; + + void * malloc_host; + cudaMallocHost((void **)&malloc_host, size); + // CHECK: dpct::pointer_attributes.get( + cuPointerGetAttributes( + numAttributes, + attributes, + attributeValues, + malloc_host + ); + std::cout << "====== Malloc Host Attributes =======" << std::endl; + std::cout << "malloc host " << malloc_host << std::endl; + std::cout << deviceID << std::endl; + std::cout << memType << std::endl; + std::cout << hostPtr << std::endl; + std::cout << devPtr << std::endl; + std::cout << isManaged << std::endl; + + // CHECK: dpct::pointer_attributes.get( + cuPointerGetAttributes( + numAttributes, + attributes, + attributeValues, + d_A + ); + std::cout << "====== Device Attributes =======" << std::endl; + std::cout << *static_cast(attributeValues[0]) << std::endl; + std::cout << attributeValues[1] << std::endl; + std::cout << attributeValues[2] << std::endl; + std::cout << *static_cast(attributeValues[3]) << std::endl; + std::cout << *static_cast(attributeValues[4]) << std::endl; + // CHECK: if (memType == sycl::usm::alloc::host) { + // CHECK: } else if (memType == sycl::usm::alloc::device) { + // CHECK: } else if (isManaged) { + if (memType == cudaMemoryTypeHost) { + return 0; + } else if (memType == cudaMemoryTypeDevice) { + return 1; + } else if (isManaged) { + return 2; + } + // CHECK: if (memType == sycl::usm::alloc::unknown) { + // CHECK: } else if (memType == sycl::usm::alloc::host) { + // CHECK: } else if (memType == sycl::usm::alloc::device) { + // CHECK: } else if (memType == sycl::usm::alloc::shared) { + if (memType == 0) { + return 0; + } else if (memType == 1) { + return 1; + } else if (memType == 2) { + return 2; + } else if (memType == 3) { + return 3; + } +} diff --git a/clang/test/dpct/query_api_mapping/Driver/test.cu b/clang/test/dpct/query_api_mapping/Driver/test.cu index 6a2295739fe3..14da18003670 100644 --- a/clang/test/dpct/query_api_mapping/Driver/test.cu +++ b/clang/test/dpct/query_api_mapping/Driver/test.cu @@ -253,6 +253,17 @@ // CUMEMPREFETCHASYNC-NEXT: dpct::queue_ptr cs; // CUMEMPREFETCHASYNC-NEXT: cs->prefetch(pd, s); +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuPointerGetAttributes | FileCheck %s -check-prefix=CUPOINTERGETATTRIBUTES +// CUPOINTERGETATTRIBUTES: CUDA API: +// CUPOINTERGETATTRIBUTES-NEXT: CUdeviceptr *ptr; +// CUPOINTERGETATTRIBUTES-NEXT: cuPointerGetAttributes(numAttributes /*unsigned int*/, +// CUPOINTERGETATTRIBUTES-NEXT: attributes /*CUpointer_attribute **/); +// CUPOINTERGETATTRIBUTES-NEXT: data /*void **/); +// CUPOINTERGETATTRIBUTES-NEXT: ptr /*CUdeviceptr*/); +// CUPOINTERGETATTRIBUTES-NEXT: Is migrated to: +// CUPOINTERGETATTRIBUTES-NEXT: dpct::device_ptr ptr; +// CUPOINTERGETATTRIBUTES-NEXT: dpct::pointer_attributes.get(numAttributes, attributes, data, ptr); + /// Stream Management // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuStreamAddCallback | FileCheck %s -check-prefix=CUSTREAMADDCALLBACK diff --git a/clang/test/dpct/query_api_mapping/test_all.cu b/clang/test/dpct/query_api_mapping/test_all.cu index 3ef020a28824..6b3b688dbe98 100644 --- a/clang/test/dpct/query_api_mapping/test_all.cu +++ b/clang/test/dpct/query_api_mapping/test_all.cu @@ -693,6 +693,7 @@ // CHECK-NEXT: cuModuleLoadDataEx // CHECK-NEXT: cuModuleUnload // CHECK-NEXT: cuOccupancyMaxActiveBlocksPerMultiprocessor +// CHECK-NEXT: cuPointerGetAttributes // CHECK-NEXT: cuStreamAddCallback // CHECK-NEXT: cuStreamAttachMemAsync // CHECK-NEXT: cuStreamCreate diff --git a/docs/dev_guide/api-mapping-status/Runtime_and_Driver_API_migration_status.csv b/docs/dev_guide/api-mapping-status/Runtime_and_Driver_API_migration_status.csv index d86c879e3c46..256e5a3025d4 100644 --- a/docs/dev_guide/api-mapping-status/Runtime_and_Driver_API_migration_status.csv +++ b/docs/dev_guide/api-mapping-status/Runtime_and_Driver_API_migration_status.csv @@ -1577,7 +1577,7 @@ cuMemPrefetchAsync,YES, Partial cuMemRangeGetAttribute,NO, cuMemRangeGetAttributes,NO, cuPointerGetAttribute,NO, -cuPointerGetAttributes,NO, +cuPointerGetAttributes,YES, Partial cuPointerSetAttribute,NO, cuGraphicsVDPAURegisterOutputSurface,NO, cuGraphicsVDPAURegisterVideoSurface,NO, From 9c1d7c138306c95e1dae58f0774892a34669383b Mon Sep 17 00:00:00 2001 From: Teja Alaghari Date: Fri, 12 Jul 2024 10:44:40 +0800 Subject: [PATCH 02/10] Corrected format and test case --- clang/examples/DPCT/Driver/cuPointerGetAttributes.cu | 1 - clang/runtime/dpct-rt/include/dpct/memory.hpp | 5 ++--- clang/test/dpct/query_api_mapping/Driver/test.cu | 3 +-- 3 files changed, 3 insertions(+), 6 deletions(-) diff --git a/clang/examples/DPCT/Driver/cuPointerGetAttributes.cu b/clang/examples/DPCT/Driver/cuPointerGetAttributes.cu index 921fc777d1d0..3c611a5da17a 100644 --- a/clang/examples/DPCT/Driver/cuPointerGetAttributes.cu +++ b/clang/examples/DPCT/Driver/cuPointerGetAttributes.cu @@ -6,4 +6,3 @@ void test(unsigned int numAttributes, CUpointer_attribute *attributes, ptr /*CUdeviceptr*/); // End } - diff --git a/clang/runtime/dpct-rt/include/dpct/memory.hpp b/clang/runtime/dpct-rt/include/dpct/memory.hpp index 94a33e77b717..c3c258cb9cb3 100644 --- a/clang/runtime/dpct-rt/include/dpct/memory.hpp +++ b/clang/runtime/dpct-rt/include/dpct/memory.hpp @@ -1693,9 +1693,8 @@ class pointer_attributes { #endif } - static void get(unsigned int numAttributes, - attribute_type *attributes, void **data, - device_ptr ptr) { + static void get(unsigned int numAttributes, attribute_type *attributes, + void **data, device_ptr ptr) { pointer_attributes sycl_attributes; sycl_attributes.init(ptr); diff --git a/clang/test/dpct/query_api_mapping/Driver/test.cu b/clang/test/dpct/query_api_mapping/Driver/test.cu index 14da18003670..c6796c1be48e 100644 --- a/clang/test/dpct/query_api_mapping/Driver/test.cu +++ b/clang/test/dpct/query_api_mapping/Driver/test.cu @@ -257,8 +257,7 @@ // CUPOINTERGETATTRIBUTES: CUDA API: // CUPOINTERGETATTRIBUTES-NEXT: CUdeviceptr *ptr; // CUPOINTERGETATTRIBUTES-NEXT: cuPointerGetAttributes(numAttributes /*unsigned int*/, -// CUPOINTERGETATTRIBUTES-NEXT: attributes /*CUpointer_attribute **/); -// CUPOINTERGETATTRIBUTES-NEXT: data /*void **/); +// CUPOINTERGETATTRIBUTES-NEXT: attributes /*CUpointer_attribute **/, data /*void **/, // CUPOINTERGETATTRIBUTES-NEXT: ptr /*CUdeviceptr*/); // CUPOINTERGETATTRIBUTES-NEXT: Is migrated to: // CUPOINTERGETATTRIBUTES-NEXT: dpct::device_ptr ptr; From 1c985996a2377971dfac79b8ee6218db0d1794fb Mon Sep 17 00:00:00 2001 From: Teja Alaghari Date: Tue, 16 Jul 2024 19:09:48 +0800 Subject: [PATCH 03/10] Corrected test case --- clang/examples/DPCT/Driver/cuPointerGetAttributes.cu | 2 +- clang/test/dpct/query_api_mapping/Driver/test.cu | 4 +--- 2 files changed, 2 insertions(+), 4 deletions(-) diff --git a/clang/examples/DPCT/Driver/cuPointerGetAttributes.cu b/clang/examples/DPCT/Driver/cuPointerGetAttributes.cu index 3c611a5da17a..138b62dee941 100644 --- a/clang/examples/DPCT/Driver/cuPointerGetAttributes.cu +++ b/clang/examples/DPCT/Driver/cuPointerGetAttributes.cu @@ -1,7 +1,7 @@ void test(unsigned int numAttributes, CUpointer_attribute *attributes, void **data, CUdeviceptr ptr) { // Start - cuPointerGetAttribute(numAttributes /*unsigned int*/, + cuPointerGetAttributes(numAttributes /*unsigned int*/, attributes /*CUpointer_attribute **/, data /*void ***/, ptr /*CUdeviceptr*/); // End diff --git a/clang/test/dpct/query_api_mapping/Driver/test.cu b/clang/test/dpct/query_api_mapping/Driver/test.cu index c6796c1be48e..6677cda9d562 100644 --- a/clang/test/dpct/query_api_mapping/Driver/test.cu +++ b/clang/test/dpct/query_api_mapping/Driver/test.cu @@ -255,12 +255,10 @@ // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuPointerGetAttributes | FileCheck %s -check-prefix=CUPOINTERGETATTRIBUTES // CUPOINTERGETATTRIBUTES: CUDA API: -// CUPOINTERGETATTRIBUTES-NEXT: CUdeviceptr *ptr; // CUPOINTERGETATTRIBUTES-NEXT: cuPointerGetAttributes(numAttributes /*unsigned int*/, -// CUPOINTERGETATTRIBUTES-NEXT: attributes /*CUpointer_attribute **/, data /*void **/, +// CUPOINTERGETATTRIBUTES-NEXT: attributes /*CUpointer_attribute **/, data /*void ***/, // CUPOINTERGETATTRIBUTES-NEXT: ptr /*CUdeviceptr*/); // CUPOINTERGETATTRIBUTES-NEXT: Is migrated to: -// CUPOINTERGETATTRIBUTES-NEXT: dpct::device_ptr ptr; // CUPOINTERGETATTRIBUTES-NEXT: dpct::pointer_attributes.get(numAttributes, attributes, data, ptr); /// Stream Management From 1da3c0ea067ee02bece2d297354d49b548e1773f Mon Sep 17 00:00:00 2001 From: Teja Alaghari Date: Wed, 17 Jul 2024 16:29:21 +0800 Subject: [PATCH 04/10] Adjusted the test case --- clang/lib/DPCT/ASTTraversal.cpp | 3 +- clang/lib/DPCT/MapNames.cpp | 11 +++++ clang/test/dpct/cuPointerGetAttributes.cu | 57 +++++++++-------------- 3 files changed, 36 insertions(+), 35 deletions(-) diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index 3e73f2b51919..4ec42e4c8b45 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -3218,7 +3218,8 @@ void EnumConstantRule::registerMatcher(MatchFinder &MF) { "libraryPropertyType_t", "cudaDataType_t", "cublasComputeType_t", "CUmem_advise_enum", "cufftType_t", "cufftType", "cudaMemoryType", - "CUctx_flags_enum", "CUpointer_attribute_enum"))), + "CUctx_flags_enum", "CUpointer_attribute_enum", + "CUmemorytype_enum"))), matchesName("CUDNN_.*"), matchesName("CUSOLVER_.*"))))) .bind("EnumConstant"), this); diff --git a/clang/lib/DPCT/MapNames.cpp b/clang/lib/DPCT/MapNames.cpp index fd8879215e09..450c9d46157c 100644 --- a/clang/lib/DPCT/MapNames.cpp +++ b/clang/lib/DPCT/MapNames.cpp @@ -1221,6 +1221,17 @@ void MapNames::setExplicitNamespaceMap() { std::make_shared(getDpctNamespace() + "attribute_type::allowed_handle_types")}, + // enum CUmemorytype Type + {"CU_MEMORYTYPE_HOST", + std::make_shared(getClNamespace() + "usm::alloc::host", + HelperFeatureEnum::device_ext)}, + {"CU_MEMORYTYPE_DEVICE", + std::make_shared(getClNamespace() + "usm::alloc::device", + HelperFeatureEnum::device_ext)}, + {"CU_MEMORYTYPE_UNIFIED", + std::make_shared(getClNamespace() + "usm::alloc::shared", + HelperFeatureEnum::device_ext)}, + // enum CUlimit {"CU_LIMIT_PRINTF_FIFO_SIZE", std::make_shared("INT_MAX")}, diff --git a/clang/test/dpct/cuPointerGetAttributes.cu b/clang/test/dpct/cuPointerGetAttributes.cu index 309ef958ade0..8ae3de32a2ed 100644 --- a/clang/test/dpct/cuPointerGetAttributes.cu +++ b/clang/test/dpct/cuPointerGetAttributes.cu @@ -1,12 +1,13 @@ -// UNSUPPORTED: v11.0, v11.1, v11.2, v11.3, v11.4, v11.5, v11.6, v11.7, v11.8, v12.0, v12.1, v12.2, v12.3 , v12.4 -// UNSUPPORTED: cuda-11.0, cuda-11.1, cuda-11.2, cuda-11.3, cuda-11.4, cuda-11.5, cuda-11.6, cuda-11.7, cuda-11.8, cuda-12.0, cuda-12.1, cuda-12.2, cuda-12.3, cuda-12.4 -// RUN: dpct --format-range=none -out-root %T/Out/cudaPointerAttributes %s --cuda-include-path="%cuda-path/include" -// RUN: FileCheck %s --match-full-lines --input-file %T/Out/cudaPointerAttributes/cudaPointerAttributes.dp.cpp -// RUN: %if build_lit %{icpx -c -fsycl %T/Out/cudaPointerAttributes/cudaPointerAttributes.dp.cpp -o %T/Out/cudaPointerAttributes/cudaPointerAttributes.dp.o %} +// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2 +// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2 +// RUN: dpct --format-range=none -out-root %T/Out/cuPointerGetAttributes %s --cuda-include-path="%cuda-path/include" +// RUN: FileCheck %s --match-full-lines --input-file %T/Out/cuPointerGetAttributes/cuPointerGetAttributes.dp.cpp +// RUN: %if build_lit %{icpx -c -fsycl %T/Out/cuPointerGetAttributes/cuPointerGetAttributes.dp.cpp -o %T/Out/cuPointerGetAttributes/cuPointerGetAttributes.dp.o %} #include #include #include #include + int main() { int N = 2048; size_t size = N * sizeof(float); @@ -17,14 +18,15 @@ int main() { cudaMalloc((void **)&d_A, size); cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); - cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice); unsigned int numAttributes = 5; // CHECK: dpct::attribute_type attributes[] = { - // CHECK: dpct::attribute_type::memory_type, dpct::attribute_type::device_pointer, - // CHECK: dpct::attribute_type::host_pointer, dpct::attribute_type::is_managed, - // CHECK: dpct::attribute_type::device_id}; + // CHECK: dpct::attribute_type::memory_type, + // CHECK: dpct::attribute_type::device_pointer, + // CHECK: dpct::attribute_type::host_pointer, + // CHECK: dpct::attribute_type::is_managed, + // CHECK: dpct::attribute_type::device_id CUpointer_attribute attributes[] = { CU_POINTER_ATTRIBUTE_MEMORY_TYPE, CU_POINTER_ATTRIBUTE_DEVICE_POINTER, @@ -34,11 +36,11 @@ int main() { }; int memType; - // CHECK: dpct::device_ptr devPtr; - CUdeviceptr devPtr; void* hostPtr; unsigned int isManaged; int deviceID; + // CHECK: dpct::device_ptr devPtr; + CUdeviceptr devPtr; void* attributeValues[] = { &memType, @@ -48,12 +50,12 @@ int main() { &deviceID }; - // CHECK: dpct::pointer_attributes.get( + // CHECK: dpct::pointer_attributes.get(numAttributes, attributes, attributeValues, (dpct::device_ptr) h_A); cuPointerGetAttributes( numAttributes, attributes, attributeValues, - h_A + (CUdeviceptr) h_A ); std::cout << "====== Host Attributes =======" << std::endl; @@ -65,12 +67,12 @@ int main() { void * malloc_host; cudaMallocHost((void **)&malloc_host, size); - // CHECK: dpct::pointer_attributes.get( + // CHECK: dpct::pointer_attributes.get(numAttributes, attributes, attributeValues, (dpct::device_ptr) malloc_host); cuPointerGetAttributes( numAttributes, attributes, attributeValues, - malloc_host + (CUdeviceptr) malloc_host ); std::cout << "====== Malloc Host Attributes =======" << std::endl; std::cout << "malloc host " << malloc_host << std::endl; @@ -80,12 +82,12 @@ int main() { std::cout << devPtr << std::endl; std::cout << isManaged << std::endl; - // CHECK: dpct::pointer_attributes.get( + // CHECK: dpct::pointer_attributes.get(numAttributes, attributes, attributeValues, (dpct::device_ptr) d_A); cuPointerGetAttributes( numAttributes, attributes, attributeValues, - d_A + (CUdeviceptr) d_A ); std::cout << "====== Device Attributes =======" << std::endl; std::cout << *static_cast(attributeValues[0]) << std::endl; @@ -93,27 +95,14 @@ int main() { std::cout << attributeValues[2] << std::endl; std::cout << *static_cast(attributeValues[3]) << std::endl; std::cout << *static_cast(attributeValues[4]) << std::endl; + // CHECK: if (memType == sycl::usm::alloc::host) { - // CHECK: } else if (memType == sycl::usm::alloc::device) { - // CHECK: } else if (isManaged) { - if (memType == cudaMemoryTypeHost) { + if (memType == CU_MEMORYTYPE_HOST) { return 0; - } else if (memType == cudaMemoryTypeDevice) { - return 1; - } else if (isManaged) { - return 2; - } - // CHECK: if (memType == sycl::usm::alloc::unknown) { - // CHECK: } else if (memType == sycl::usm::alloc::host) { // CHECK: } else if (memType == sycl::usm::alloc::device) { - // CHECK: } else if (memType == sycl::usm::alloc::shared) { - if (memType == 0) { - return 0; - } else if (memType == 1) { + } else if (memType == CU_MEMORYTYPE_DEVICE) { return 1; - } else if (memType == 2) { + } else if (isManaged) { return 2; - } else if (memType == 3) { - return 3; } } From 54be268996539449198282a5c8c8589c264f5ac5 Mon Sep 17 00:00:00 2001 From: Teja Alaghari Date: Wed, 17 Jul 2024 18:03:22 +0800 Subject: [PATCH 05/10] Corrected migration logic --- clang/lib/DPCT/APINamesMemory.inc | 2 +- clang/lib/DPCT/ASTTraversal.cpp | 3 ++- clang/lib/DPCT/MapNames.cpp | 4 ++++ clang/test/dpct/cuPointerGetAttributes.cu | 9 +++++---- clang/test/dpct/query_api_mapping/Driver/test.cu | 2 +- 5 files changed, 13 insertions(+), 7 deletions(-) diff --git a/clang/lib/DPCT/APINamesMemory.inc b/clang/lib/DPCT/APINamesMemory.inc index 1350d3a0ccde..1ce02f34191c 100644 --- a/clang/lib/DPCT/APINamesMemory.inc +++ b/clang/lib/DPCT/APINamesMemory.inc @@ -402,7 +402,7 @@ ASSIGNABLE_FACTORY( FEATURE_REQUEST_FACTORY( HelperFeatureEnum::device_ext, CALL_FACTORY_ENTRY( - "cuPointerGetAttributes", CALL(MapNames::getDpctNamespace() + "pointer_attributes.get", + "cuPointerGetAttributes", CALL(MapNames::getDpctNamespace() + "pointer_attributes::get", ARG(0), ARG(1), ARG(2), ARG(3))))) ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index 4ec42e4c8b45..bb464b84b563 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -1728,7 +1728,8 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) { "cooperative_groups::__v1::thread_block", "libraryPropertyType_t", "libraryPropertyType", "cudaDataType_t", "cudaDataType", "cublasComputeType_t", "cublasAtomicsMode_t", "cublasMath_t", - "CUmem_advise_enum", "CUmem_advise", "thrust::tuple_element", + "CUmem_advise_enum", "CUmem_advise", "CUmemorytype", + "CUmemorytype_enum", "thrust::tuple_element", "thrust::tuple_size", "thrust::zip_iterator", "cudaPointerAttributes", "CUpointer_attribute", "cusolverEigRange_t", "cudaUUID_t", "cusolverDnFunction_t", diff --git a/clang/lib/DPCT/MapNames.cpp b/clang/lib/DPCT/MapNames.cpp index 450c9d46157c..232231f875bc 100644 --- a/clang/lib/DPCT/MapNames.cpp +++ b/clang/lib/DPCT/MapNames.cpp @@ -372,6 +372,10 @@ void MapNames::setExplicitNamespaceMap() { ? getClNamespace() + "ext::oneapi::experimental::queue_state" : "cudaStreamCaptureStatus")}, {"CUmem_advise", std::make_shared("int")}, + {"CUmemorytype", + std::make_shared(getClNamespace() + "usm::alloc")}, + {"CUmemorytype_enum", + std::make_shared(getClNamespace() + "usm::alloc")}, {"cudaPos", std::make_shared(getClNamespace() + "id<3>")}, {"cudaExtent", std::make_shared(getClNamespace() + "range<3>")}, diff --git a/clang/test/dpct/cuPointerGetAttributes.cu b/clang/test/dpct/cuPointerGetAttributes.cu index 8ae3de32a2ed..929ec704bb10 100644 --- a/clang/test/dpct/cuPointerGetAttributes.cu +++ b/clang/test/dpct/cuPointerGetAttributes.cu @@ -35,7 +35,8 @@ int main() { CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL }; - int memType; + // CHECK: sycl::usm::alloc memType; + CUmemorytype memType; void* hostPtr; unsigned int isManaged; int deviceID; @@ -50,7 +51,7 @@ int main() { &deviceID }; - // CHECK: dpct::pointer_attributes.get(numAttributes, attributes, attributeValues, (dpct::device_ptr) h_A); + // CHECK: dpct::pointer_attributes::get(numAttributes, attributes, attributeValues, (dpct::device_ptr) h_A); cuPointerGetAttributes( numAttributes, attributes, @@ -67,7 +68,7 @@ int main() { void * malloc_host; cudaMallocHost((void **)&malloc_host, size); - // CHECK: dpct::pointer_attributes.get(numAttributes, attributes, attributeValues, (dpct::device_ptr) malloc_host); + // CHECK: dpct::pointer_attributes::get(numAttributes, attributes, attributeValues, (dpct::device_ptr) malloc_host); cuPointerGetAttributes( numAttributes, attributes, @@ -82,7 +83,7 @@ int main() { std::cout << devPtr << std::endl; std::cout << isManaged << std::endl; - // CHECK: dpct::pointer_attributes.get(numAttributes, attributes, attributeValues, (dpct::device_ptr) d_A); + // CHECK: dpct::pointer_attributes::get(numAttributes, attributes, attributeValues, (dpct::device_ptr) d_A); cuPointerGetAttributes( numAttributes, attributes, diff --git a/clang/test/dpct/query_api_mapping/Driver/test.cu b/clang/test/dpct/query_api_mapping/Driver/test.cu index 6677cda9d562..b762b8b24860 100644 --- a/clang/test/dpct/query_api_mapping/Driver/test.cu +++ b/clang/test/dpct/query_api_mapping/Driver/test.cu @@ -259,7 +259,7 @@ // CUPOINTERGETATTRIBUTES-NEXT: attributes /*CUpointer_attribute **/, data /*void ***/, // CUPOINTERGETATTRIBUTES-NEXT: ptr /*CUdeviceptr*/); // CUPOINTERGETATTRIBUTES-NEXT: Is migrated to: -// CUPOINTERGETATTRIBUTES-NEXT: dpct::pointer_attributes.get(numAttributes, attributes, data, ptr); +// CUPOINTERGETATTRIBUTES-NEXT: dpct::pointer_attributes::get(numAttributes, attributes, data, ptr); /// Stream Management From 5e89557e7e6d6457bc22111652741579e94f48f6 Mon Sep 17 00:00:00 2001 From: Teja Alaghari Date: Wed, 17 Jul 2024 23:34:54 +0800 Subject: [PATCH 06/10] Corrected type cast of CUmemorytype enum --- clang/test/dpct/cuPointerGetAttributes.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/test/dpct/cuPointerGetAttributes.cu b/clang/test/dpct/cuPointerGetAttributes.cu index 929ec704bb10..9c0192f896ad 100644 --- a/clang/test/dpct/cuPointerGetAttributes.cu +++ b/clang/test/dpct/cuPointerGetAttributes.cu @@ -61,7 +61,7 @@ int main() { std::cout << "====== Host Attributes =======" << std::endl; std::cout << deviceID << std::endl; - std::cout << memType << std::endl; + std::cout << static_cast(memType) << std::endl; std::cout << hostPtr << std::endl; std::cout << devPtr << std::endl; std::cout << isManaged << std::endl; @@ -78,7 +78,7 @@ int main() { std::cout << "====== Malloc Host Attributes =======" << std::endl; std::cout << "malloc host " << malloc_host << std::endl; std::cout << deviceID << std::endl; - std::cout << memType << std::endl; + std::cout << static_cast(memType) << std::endl; std::cout << hostPtr << std::endl; std::cout << devPtr << std::endl; std::cout << isManaged << std::endl; From fc6153b4357b04676d7b9193b6155798867ccc0c Mon Sep 17 00:00:00 2001 From: Teja Alaghari Date: Fri, 19 Jul 2024 18:16:09 +0800 Subject: [PATCH 07/10] Included attribute_types class under pointer_attributes class --- clang/lib/DPCT/MapNames.cpp | 30 +++++++-------- clang/runtime/dpct-rt/include/dpct/memory.hpp | 38 ++++++++----------- clang/test/dpct/cuPointerGetAttributes.cu | 12 +++--- 3 files changed, 36 insertions(+), 44 deletions(-) diff --git a/clang/lib/DPCT/MapNames.cpp b/clang/lib/DPCT/MapNames.cpp index 232231f875bc..9ca628d95964 100644 --- a/clang/lib/DPCT/MapNames.cpp +++ b/clang/lib/DPCT/MapNames.cpp @@ -109,7 +109,7 @@ void MapNames::setExplicitNamespaceMap() { std::make_shared(getDpctNamespace() + "kernel_function", HelperFeatureEnum::device_ext)}, {"CUpointer_attribute", - std::make_shared(getDpctNamespace() + "attribute_type")}, + std::make_shared(getDpctNamespace() + "pointer_attributes::type")}, {"cudaPointerAttributes", std::make_shared(getDpctNamespace() + "pointer_attributes", HelperFeatureEnum::device_ext)}, @@ -1184,46 +1184,46 @@ void MapNames::setExplicitNamespaceMap() { // enum CUpointer_attribute {"CU_POINTER_ATTRIBUTE_CONTEXT", std::make_shared(getDpctNamespace() + - "attribute_type::context")}, + "pointer_attributes::type::unsupported")}, {"CU_POINTER_ATTRIBUTE_MEMORY_TYPE", std::make_shared(getDpctNamespace() + - "attribute_type::memory_type")}, + "pointer_attributes::type::memory_type")}, {"CU_POINTER_ATTRIBUTE_DEVICE_POINTER", std::make_shared(getDpctNamespace() + - "attribute_type::device_pointer")}, + "pointer_attributes::type::device_pointer")}, {"CU_POINTER_ATTRIBUTE_HOST_POINTER", std::make_shared(getDpctNamespace() + - "attribute_type::host_pointer")}, + "pointer_attributes::type::host_pointer")}, {"CU_POINTER_ATTRIBUTE_P2P_TOKENS", std::make_shared(getDpctNamespace() + - "attribute_type::p2p_tokens")}, + "pointer_attributes::type::unsupported")}, {"CU_POINTER_ATTRIBUTE_SYNC_MEMOPS", std::make_shared(getDpctNamespace() + - "attribute_type::sync_memps")}, + "pointer_attributes::type::unsupported")}, {"CU_POINTER_ATTRIBUTE_BUFFER_ID", std::make_shared(getDpctNamespace() + - "attribute_type::buffer_id")}, + "pointer_attributes::type::unsupported")}, {"CU_POINTER_ATTRIBUTE_IS_MANAGED", std::make_shared(getDpctNamespace() + - "attribute_type::is_managed")}, + "pointer_attributes::type::is_managed")}, {"CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL", std::make_shared(getDpctNamespace() + - "attribute_type::device_id")}, + "pointer_attributes::type::device_id")}, {"CU_POINTER_ATTRIBUTE_IS_LEGACY_CUDA_IPC_CAPABLE", std::make_shared(getDpctNamespace() + - "attribute_type::is_legacy_ipc_capable")}, + "pointer_attributes::type::unsupported")}, {"CU_POINTER_ATTRIBUTE_RANGE_START_ADDR", std::make_shared(getDpctNamespace() + - "attribute_type::range_start_addr")}, + "pointer_attributes::type::unsupported")}, {"CU_POINTER_ATTRIBUTE_RANGE_SIZE", std::make_shared(getDpctNamespace() + - "attribute_type::range_size")}, + "pointer_attributes::type::unsupported")}, {"CU_POINTER_ATTRIBUTE_MAPPED", std::make_shared(getDpctNamespace() + - "attribute_type::mapped")}, + "pointer_attributes::type::unsupported")}, {"CU_POINTER_ATTRIBUTE_ALLOWED_HANDLE_TYPES", std::make_shared(getDpctNamespace() + - "attribute_type::allowed_handle_types")}, + "pointer_attributes::type::unsupported")}, // enum CUmemorytype Type {"CU_MEMORYTYPE_HOST", diff --git a/clang/runtime/dpct-rt/include/dpct/memory.hpp b/clang/runtime/dpct-rt/include/dpct/memory.hpp index c3c258cb9cb3..5afb17d537b7 100644 --- a/clang/runtime/dpct-rt/include/dpct/memory.hpp +++ b/clang/runtime/dpct-rt/include/dpct/memory.hpp @@ -1653,25 +1653,17 @@ template using usm_device_allocator = detail::deprecated::usm_allocator; } // namespace deprecated -enum class attribute_type { - context, - memory_type, - device_pointer, - host_pointer, - p2p_tokens, - sync_memps, - buffer_id, - is_managed, - device_id, - is_legacy_ipc_capable, - range_start_addr, - range_size, - mapped, - allowed_handle_types -}; - class pointer_attributes { public: + enum class type { + memory_type, + device_pointer, + host_pointer, + is_managed, + device_id, + unsupported + }; + void init(const void *ptr, sycl::queue &q = dpct::get_default_queue()) { #ifdef DPCT_USM_LEVEL_NONE @@ -1693,7 +1685,7 @@ class pointer_attributes { #endif } - static void get(unsigned int numAttributes, attribute_type *attributes, + static void get(unsigned int numAttributes, type *attributes, void **data, device_ptr ptr) { pointer_attributes sycl_attributes; @@ -1701,23 +1693,23 @@ class pointer_attributes { for (int i = 0; i < numAttributes; i++) { switch (attributes[i]) { - case attribute_type::memory_type: + case type::memory_type: *static_cast(data[i]) = static_cast(sycl_attributes.get_memory_type()); break; - case attribute_type::device_pointer: + case type::device_pointer: *(reinterpret_cast(data[i])) = const_cast(sycl_attributes.get_device_pointer()); break; - case attribute_type::host_pointer: + case type::host_pointer: *(reinterpret_cast(data[i])) = const_cast(sycl_attributes.get_host_pointer()); break; - case attribute_type::is_managed: + case type::is_managed: *static_cast(data[i]) = sycl_attributes.is_memory_shared(); break; - case attribute_type::device_id: + case type::device_id: *static_cast(data[i]) = sycl_attributes.get_device_id(); break; default: diff --git a/clang/test/dpct/cuPointerGetAttributes.cu b/clang/test/dpct/cuPointerGetAttributes.cu index 9c0192f896ad..b9daba8f59c3 100644 --- a/clang/test/dpct/cuPointerGetAttributes.cu +++ b/clang/test/dpct/cuPointerGetAttributes.cu @@ -21,12 +21,12 @@ int main() { unsigned int numAttributes = 5; - // CHECK: dpct::attribute_type attributes[] = { - // CHECK: dpct::attribute_type::memory_type, - // CHECK: dpct::attribute_type::device_pointer, - // CHECK: dpct::attribute_type::host_pointer, - // CHECK: dpct::attribute_type::is_managed, - // CHECK: dpct::attribute_type::device_id + // CHECK: dpct::pointer_attributes::type attributes[] = { + // CHECK: dpct::pointer_attributes::type::memory_type, + // CHECK: dpct::pointer_attributes::type::device_pointer, + // CHECK: dpct::pointer_attributes::type::host_pointer, + // CHECK: dpct::pointer_attributes::type::is_managed, + // CHECK: dpct::pointer_attributes::type::device_id CUpointer_attribute attributes[] = { CU_POINTER_ATTRIBUTE_MEMORY_TYPE, CU_POINTER_ATTRIBUTE_DEVICE_POINTER, From 4af59b11221dab4b5674826415d86f79e6f571f4 Mon Sep 17 00:00:00 2001 From: Teja Alaghari Date: Fri, 2 Aug 2024 15:18:21 +0800 Subject: [PATCH 08/10] Added new test case scenario for unsupported attr --- .../DPCT/Driver/cuPointerGetAttributes.cu | 6 ++--- clang/lib/DPCT/APINames.inc | 4 ++-- clang/test/dpct/cuPointerGetAttributes.cu | 23 +++++++++++++++++++ .../dpct/query_api_mapping/Driver/test.cu | 6 ++--- 4 files changed, 31 insertions(+), 8 deletions(-) diff --git a/clang/examples/DPCT/Driver/cuPointerGetAttributes.cu b/clang/examples/DPCT/Driver/cuPointerGetAttributes.cu index 138b62dee941..2535eb34bf77 100644 --- a/clang/examples/DPCT/Driver/cuPointerGetAttributes.cu +++ b/clang/examples/DPCT/Driver/cuPointerGetAttributes.cu @@ -1,8 +1,8 @@ -void test(unsigned int numAttributes, CUpointer_attribute *attributes, +void test(unsigned int numAttr, CUpointer_attribute *attr, void **data, CUdeviceptr ptr) { // Start - cuPointerGetAttributes(numAttributes /*unsigned int*/, - attributes /*CUpointer_attribute **/, data /*void ***/, + cuPointerGetAttributes(numAttr /*unsigned int*/, + attr /*CUpointer_attribute **/, data /*void ***/, ptr /*CUdeviceptr*/); // End } diff --git a/clang/lib/DPCT/APINames.inc b/clang/lib/DPCT/APINames.inc index 1d57141e8f9c..20b4111a14cc 100644 --- a/clang/lib/DPCT/APINames.inc +++ b/clang/lib/DPCT/APINames.inc @@ -251,7 +251,7 @@ ENTRY(cudaMemPoolSetAttribute, cudaMemPoolSetAttribute, false, NO_FLAG, P4, "com ENTRY(cudaMemPoolTrimTo, cudaMemPoolTrimTo, false, NO_FLAG, P4, "comment") // unified addressing functions of runtime API -ENTRY(cudaPointerGetAttributes, cudaPointerGetAttributes, true, NO_FLAG, P0, "Successful") +ENTRY(cudaPointerGetAttributes, cudaPointerGetAttributes, true, NO_FLAG, P0, "Partial: Only 5 attributes are supported") // peer device memory access functions of runtime API ENTRY(cudaDeviceCanAccessPeer, cudaDeviceCanAccessPeer, true, NO_FLAG, P0, "DPCT1031") @@ -1763,7 +1763,7 @@ ENTRY(cuMemPrefetchAsync, cuMemPrefetchAsync, true, NO_FLAG, P4, "Partial: USM o ENTRY(cuMemPrefetchAsync_v2, cuMemPrefetchAsync_v2, false, NO_FLAG, P4, "Partial: USM only, DPCT1007") ENTRY(cuMemRangeGetAttribute, cuMemRangeGetAttribute, false, NO_FLAG, P4, "comment") ENTRY(cuMemRangeGetAttributes, cuMemRangeGetAttributes, false, NO_FLAG, P4, "comment") -ENTRY(cuPointerGetAttribute, cuPointerGetAttribute, false, NO_FLAG, P4, "comment") +ENTRY(cuPointerGetAttribute, cuPointerGetAttribute, true, NO_FLAG, P4, "Successful") ENTRY(cuPointerGetAttributes, cuPointerGetAttributes, false, NO_FLAG, P4, "comment") ENTRY(cuPointerSetAttribute, cuPointerSetAttribute, false, NO_FLAG, P4, "comment") diff --git a/clang/test/dpct/cuPointerGetAttributes.cu b/clang/test/dpct/cuPointerGetAttributes.cu index b9daba8f59c3..e9956b5aa54d 100644 --- a/clang/test/dpct/cuPointerGetAttributes.cu +++ b/clang/test/dpct/cuPointerGetAttributes.cu @@ -35,6 +35,14 @@ int main() { CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL }; + // CHECK: dpct::pointer_attributes::type attributes_unsupported[] = { + // CHECK: dpct::pointer_attributes::type::unsupported, + // CHECK: dpct::pointer_attributes::type::unsupported + CUpointer_attribute attributes_unsupported[] = { + CU_POINTER_ATTRIBUTE_CONTEXT, + CU_POINTER_ATTRIBUTE_IS_LEGACY_CUDA_IPC_CAPABLE + }; + // CHECK: sycl::usm::alloc memType; CUmemorytype memType; void* hostPtr; @@ -42,6 +50,8 @@ int main() { int deviceID; // CHECK: dpct::device_ptr devPtr; CUdeviceptr devPtr; + CUcontext cuCtx; + bool is_legacy_cuda_ipc_capable; void* attributeValues[] = { &memType, @@ -51,6 +61,11 @@ int main() { &deviceID }; + void* attributeValues_unsupported[] = { + &cuCtx, + &is_legacy_cuda_ipc_capable + }; + // CHECK: dpct::pointer_attributes::get(numAttributes, attributes, attributeValues, (dpct::device_ptr) h_A); cuPointerGetAttributes( numAttributes, @@ -106,4 +121,12 @@ int main() { } else if (isManaged) { return 2; } + + // CHECK: dpct::pointer_attributes::get(2, attributes_unsupported, attributeValues_unsupported, (dpct::device_ptr) d_A); + cuPointerGetAttributes( + 2, + attributes_unsupported, + attributeValues_unsupported, + (CUdeviceptr) d_A + ); } diff --git a/clang/test/dpct/query_api_mapping/Driver/test.cu b/clang/test/dpct/query_api_mapping/Driver/test.cu index b762b8b24860..1876a51b95a2 100644 --- a/clang/test/dpct/query_api_mapping/Driver/test.cu +++ b/clang/test/dpct/query_api_mapping/Driver/test.cu @@ -255,11 +255,11 @@ // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuPointerGetAttributes | FileCheck %s -check-prefix=CUPOINTERGETATTRIBUTES // CUPOINTERGETATTRIBUTES: CUDA API: -// CUPOINTERGETATTRIBUTES-NEXT: cuPointerGetAttributes(numAttributes /*unsigned int*/, -// CUPOINTERGETATTRIBUTES-NEXT: attributes /*CUpointer_attribute **/, data /*void ***/, +// CUPOINTERGETATTRIBUTES-NEXT: cuPointerGetAttributes(numAttr /*unsigned int*/, +// CUPOINTERGETATTRIBUTES-NEXT: attr /*CUpointer_attribute **/, data /*void ***/, // CUPOINTERGETATTRIBUTES-NEXT: ptr /*CUdeviceptr*/); // CUPOINTERGETATTRIBUTES-NEXT: Is migrated to: -// CUPOINTERGETATTRIBUTES-NEXT: dpct::pointer_attributes::get(numAttributes, attributes, data, ptr); +// CUPOINTERGETATTRIBUTES-NEXT: dpct::pointer_attributes::get(numAttr, attr, data, ptr); /// Stream Management From 82d700ffa667f0b091ba8af7d61d98475df3be1f Mon Sep 17 00:00:00 2001 From: Teja Alaghari <98731765+TejaX-Alaghari@users.noreply.github.com> Date: Fri, 23 Aug 2024 12:22:36 +0530 Subject: [PATCH 09/10] Update APINames.inc to include correct status --- clang/lib/DPCT/APINames.inc | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/clang/lib/DPCT/APINames.inc b/clang/lib/DPCT/APINames.inc index 7febe0fb4fcd..20898e288cb0 100644 --- a/clang/lib/DPCT/APINames.inc +++ b/clang/lib/DPCT/APINames.inc @@ -251,7 +251,7 @@ ENTRY(cudaMemPoolSetAttribute, cudaMemPoolSetAttribute, false, NO_FLAG, P4, "com ENTRY(cudaMemPoolTrimTo, cudaMemPoolTrimTo, false, NO_FLAG, P4, "comment") // unified addressing functions of runtime API -ENTRY(cudaPointerGetAttributes, cudaPointerGetAttributes, true, NO_FLAG, P0, "Partial: Only 5 attributes are supported") +ENTRY(cudaPointerGetAttributes, cudaPointerGetAttributes, true, NO_FLAG, P0, "Successful") // peer device memory access functions of runtime API ENTRY(cudaDeviceCanAccessPeer, cudaDeviceCanAccessPeer, true, NO_FLAG, P0, "DPCT1031") @@ -1763,8 +1763,8 @@ ENTRY(cuMemPrefetchAsync, cuMemPrefetchAsync, true, NO_FLAG, P4, "Partial: USM o ENTRY(cuMemPrefetchAsync_v2, cuMemPrefetchAsync_v2, false, NO_FLAG, P4, "Partial: USM only, DPCT1007") ENTRY(cuMemRangeGetAttribute, cuMemRangeGetAttribute, false, NO_FLAG, P4, "comment") ENTRY(cuMemRangeGetAttributes, cuMemRangeGetAttributes, false, NO_FLAG, P4, "comment") -ENTRY(cuPointerGetAttribute, cuPointerGetAttribute, true, NO_FLAG, P4, "Successful") -ENTRY(cuPointerGetAttributes, cuPointerGetAttributes, false, NO_FLAG, P4, "comment") +ENTRY(cuPointerGetAttribute, cuPointerGetAttribute, false, NO_FLAG, P4, "comment") +ENTRY(cuPointerGetAttributes, cuPointerGetAttributes, true, NO_FLAG, P4, "Partial: Only 5 attributes are supported") ENTRY(cuPointerSetAttribute, cuPointerSetAttribute, false, NO_FLAG, P4, "comment") // Stream Management @@ -2304,4 +2304,4 @@ ENTRY(make_cuComplex, make_cuComplex, true, NO_FLAG, P4, "comment") ENTRY(make_cuDoubleComplex, make_cuDoubleComplex, true, NO_FLAG, P4, "comment") ENTRY(make_cuFloatComplex, make_cuFloatComplex, true, NO_FLAG, P4, "comment") ENTRY(__assert_fail, __assert_fail, true, NO_FLAG, P4, "Successful") -ENTRY(__assertfail, __assertfail, true, NO_FLAG, P4, "Successful") \ No newline at end of file +ENTRY(__assertfail, __assertfail, true, NO_FLAG, P4, "Successful") From 8f49a419e7f751d0a651ed8615d87d8c996b7b80 Mon Sep 17 00:00:00 2001 From: Teja Alaghari <98731765+TejaX-Alaghari@users.noreply.github.com> Date: Fri, 23 Aug 2024 12:30:35 +0530 Subject: [PATCH 10/10] Update memory.hpp --- clang/runtime/dpct-rt/include/dpct/memory.hpp | 1 + 1 file changed, 1 insertion(+) diff --git a/clang/runtime/dpct-rt/include/dpct/memory.hpp b/clang/runtime/dpct-rt/include/dpct/memory.hpp index 5afb17d537b7..9e885211b62e 100644 --- a/clang/runtime/dpct-rt/include/dpct/memory.hpp +++ b/clang/runtime/dpct-rt/include/dpct/memory.hpp @@ -1685,6 +1685,7 @@ class pointer_attributes { #endif } + // Query pointer propreties listed in attributes and store the results in data array static void get(unsigned int numAttributes, type *attributes, void **data, device_ptr ptr) { pointer_attributes sycl_attributes;