[Query API Mapping] Added QAM for **4 Driver (Blender) & 23 nvSHMEM APIs**

clang/examples/DPCT/Driver/cuDeviceGetP2PAttribute.cu

@@ -0,0 +1,11 @@
+void test(int value, CUdevice srcDevice, CUdevice dstDevice) {
+  // Start
+  CUdevice_P2PAttribute attrib;
+  /* 1 */ cuDeviceGetP2PAttribute(
+      &value /*int **/, CU_DEVICE_P2P_ATTRIBUTE_ACCESS_SUPPORTED,
+      srcDevice /*CUdevice*/, dstDevice /*CUdevice*/);
+  /* 2 */ cuDeviceGetP2PAttribute(
+      &value /*int **/, CU_DEVICE_P2P_ATTRIBUTE_NATIVE_ATOMIC_SUPPORTED,
+      srcDevice /*CUdevice*/, dstDevice /*CUdevice*/);
+  // End
+}

clang/examples/DPCT/Driver/cuDevicePrimaryCtxGetState.cu

@@ -0,0 +1,6 @@
+void test(CUdevice device, unsigned int flags, int active) {
+  // Start
+  cuDevicePrimaryCtxGetState(device /*CUdevice*/, &flags /*unsigned int **/,
+                             &active /*int **/);
+  // End
+}

clang/examples/DPCT/Driver/cuDevicePrimaryCtxSetFlags.cu

@@ -0,0 +1,5 @@
+void test(CUdevice device, unsigned int flags) {
+  // Start
+  cuDevicePrimaryCtxSetFlags(device /*CUdevice*/, flags /*unsigned int*/);
+  // End
+}

clang/examples/DPCT/Driver/cuMemcpy2DUnaligned.cu

@@ -0,0 +1,5 @@
+void test(const CUDA_MEMCPY2D *copy2d_desc) {
+  // Start
+  cuMemcpy2DUnaligned(copy2d_desc /*const CUDA_MEMCPY2D **/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_align.cu

@@ -0,0 +1,7 @@
+#include <nvshmem.h>
+
+void test(size_t alignment, size_t size) {
+  // Start
+  nvshmem_align(alignment /*size_t*/, size /*size_t*/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_calloc.cu

@@ -0,0 +1,7 @@
+#include <nvshmem.h>
+
+void test(size_t count, size_t size) {
+  // Start
+  nvshmem_calloc(count /*size_t*/, size /*size_t*/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_finalize.cu

@@ -0,0 +1,7 @@
+#include <nvshmem.h>
+
+void test() {
+  // Start
+  nvshmem_finalize(/*void*/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_free.cu

@@ -0,0 +1,7 @@
+#include <nvshmem.h>
+
+void test(void *ptr) {
+  // Start
+  nvshmem_free(ptr /*void **/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_info_get_name.cu

@@ -0,0 +1,7 @@
+#include <nvshmem.h>
+
+void test(char *name) {
+  // Start
+  nvshmem_info_get_name(name /*char **/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_info_get_version.cu

@@ -0,0 +1,7 @@
+#include <nvshmem.h>
+
+void test(int *major, int *minor) {
+  // Start
+  nvshmem_info_get_version(major /*int **/, minor /*int **/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_init.cu

@@ -0,0 +1,7 @@
+#include <nvshmem.h>
+
+void test() {
+  // Start
+  nvshmem_init(/*void*/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_malloc.cu

@@ -0,0 +1,7 @@
+#include <nvshmem.h>
+
+void test(size_t size) {
+  // Start
+  nvshmem_malloc(size /*size_t*/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_my_pe.cu

@@ -0,0 +1,7 @@
+#include <nvshmem.h>
+
+void test() {
+  // Start
+  nvshmem_my_pe(/*void*/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_n_pes.cu

@@ -0,0 +1,7 @@
+#include <nvshmem.h>
+
+void test() {
+  // Start
+  nvshmem_n_pes(/*void*/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_ptr.cu

@@ -0,0 +1,7 @@
+#include <nvshmem.h>
+
+void test(void *local_ptr, const void *dest, int pe) {
+  // Start
+  local_ptr = nvshmem_ptr(dest /*const void **/, pe /*int*/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_putmem_nbi.cu

@@ -0,0 +1,9 @@
+#include <nvshmem.h>
+
+__device__ __host__ void test(void *dest, const void *source, size_t nelems,
+                              int pe) {
+  // Start
+  nvshmem_putmem_nbi(dest /*void **/, source /*const void **/,
+                     nelems /*size_t*/, pe /*int*/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_putmem_signal_nbi.cu

@@ -0,0 +1,15 @@
+#include <nvshmem.h>
+
+__device__ __host__ void test(void *dest, const void *source, size_t nelems,
+                              uint64_t *sig_addr, uint64_t signal, int pe) {
+  // Start
+  /* 1 */ nvshmem_putmem_signal_nbi(dest /*void **/, source /*const void **/,
+                                    nelems /*size_t*/, sig_addr /*uint64_t **/,
+                                    signal /*uint64_t*/, NVSHMEM_SIGNAL_SET,
+                                    pe /*int*/);
+  /* 2 */ nvshmem_putmem_signal_nbi(dest /*void **/, source /*const void **/,
+                                    nelems /*size_t*/, sig_addr /*uint64_t **/,
+                                    signal /*uint64_t*/, NVSHMEM_SIGNAL_ADD,
+                                    pe /*int*/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_signal_wait_until.cu

@@ -0,0 +1,18 @@
+#include <nvshmem.h>
+
+__device__ __host__ void test(uint64_t *sig_addr, uint64_t cmp_val) {
+  // Start
+  /* 1 */ nvshmem_signal_wait_until(sig_addr /*uint64_t **/, NVSHMEM_CMP_EQ,
+                                    cmp_val /*uint64_t*/);
+  /* 2 */ nvshmem_signal_wait_until(sig_addr /*uint64_t **/, NVSHMEM_CMP_NE,
+                                    cmp_val /*uint64_t*/);
+  /* 3 */ nvshmem_signal_wait_until(sig_addr /*uint64_t **/, NVSHMEM_CMP_GT,
+                                    cmp_val /*uint64_t*/);
+  /* 4 */ nvshmem_signal_wait_until(sig_addr /*uint64_t **/, NVSHMEM_CMP_GE,
+                                    cmp_val /*uint64_t*/);
+  /* 5 */ nvshmem_signal_wait_until(sig_addr /*uint64_t **/, NVSHMEM_CMP_LT,
+                                    cmp_val /*uint64_t*/);
+  /* 6 */ nvshmem_signal_wait_until(sig_addr /*uint64_t **/, NVSHMEM_CMP_LE,
+                                    cmp_val /*uint64_t*/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_team_destroy.cu

@@ -0,0 +1,8 @@
+#include <nvshmem.h>
+
+__host__ void test() {
+  // Start
+  nvshmem_team_t team;
+  nvshmem_team_destroy(team /*nvshmem_team_t*/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_team_get_config.cu

@@ -0,0 +1,9 @@
+#include <nvshmem.h>
+
+__host__ void test(nvshmem_team_t team) {
+  // Start
+  nvshmem_team_config_t *config;
+  nvshmem_team_get_config(team /*nvshmem_team_t*/,
+                          config /*nvshmem_team_config_t **/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_team_my_pe.cu

@@ -0,0 +1,7 @@
+#include <nvshmem.h>
+
+__host__ void test(nvshmem_team_t team) {
+  // Start
+  nvshmem_team_my_pe(team /*nvshmem_team_t*/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_team_n_pes.cu

@@ -0,0 +1,7 @@
+#include <nvshmem.h>
+
+__host__ void test(nvshmem_team_t team) {
+  // Start
+  nvshmem_team_n_pes(team /*nvshmem_team_t*/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_team_split_2d.cu

@@ -0,0 +1,15 @@
+#include <nvshmem.h>
+
+__host__ void test(nvshmem_team_t parent_team, int xrange,
+                   const nvshmem_team_config_t *xaxis_config, long xaxis_mask,
+                   nvshmem_team_t *xaxis_team,
+                   const nvshmem_team_config_t *yaxis_config, long yaxis_mask,
+                   nvshmem_team_t *yaxis_team) {
+  // Start
+  nvshmem_team_split_2d(parent_team /*nvshmem_team_t*/, xrange /*int*/,
+                        xaxis_config /*const nvshmem_team_config_t **/,
+                        xaxis_mask /*long*/, xaxis_team /*nvshmem_team_t **/,
+                        yaxis_config /*const nvshmem_team_config_t **/,
+                        yaxis_mask /*long*/, yaxis_team /*nvshmem_team_t **/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_team_split_strided.cu

@@ -0,0 +1,12 @@
+#include <nvshmem.h>
+
+__host__ void test(nvshmem_team_t parent_team, int start, int stride, int size,
+                   const nvshmem_team_config_t *config, long config_mask,
+                   nvshmem_team_t *new_team) {
+  // Start
+  nvshmem_team_split_strided(
+      parent_team /*nvshmem_team_t*/, start /*int*/, stride /*int*/,
+      size /*int*/, config /*const nvshmem_team_config_t **/,
+      config_mask /*long*/, new_team /*nvshmem_team_t **/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmem_team_translate_pe.cu

@@ -0,0 +1,9 @@
+#include <nvshmem.h>
+
+__host__ void test(nvshmem_team_t src_team, int src_pe,
+                   nvshmem_team_t dest_team) {
+  // Start
+  nvshmem_team_translate_pe(src_team /*nvshmem_team_t*/, src_pe /*int*/,
+                            dest_team /*nvshmem_team_t*/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmemx_init_attr.cu

@@ -0,0 +1,11 @@
+#include <nvshmemx.h>
+
+__host__ void test() {
+  // Start
+  nvshmemx_init_attr_t attributes;
+  /* 1 */ nvshmemx_init_attr(NVSHMEMX_INIT_WITH_MPI_COMM,
+                             &attributes /*nvshmemx_init_attr_t **/);
+  /* 2 */ nvshmemx_init_attr(NVSHMEMX_INIT_WITH_SHMEM,
+                             &attributes /*nvshmemx_init_attr_t **/);
+  // End
+}

clang/examples/DPCT/nvSHMEM/nvshmemx_signal_op.cu

@@ -0,0 +1,10 @@
+#include <nvshmemx.h>
+
+__device__ __host__ void test(uint64_t *sig_addr, uint64_t signal, int pe) {
+  // Start
+  /* 1 */ nvshmemx_signal_op(sig_addr /*uint64_t **/, signal /*uint64_t*/,
+                             NVSHMEM_SIGNAL_SET, pe /*int*/);
+  /* 2 */ nvshmemx_signal_op(sig_addr /*uint64_t **/, signal /*uint64_t*/,
+                             NVSHMEM_SIGNAL_ADD, pe /*int*/);
+  // End
+}

clang/test/dpct/query_api_mapping/Driver/test.cu

@@ -90,6 +90,20 @@
 // CUDEVICEGETATTRIBUTE-NEXT:  /* 15 */ *pi = dpct::get_device(d).get_minor_version();
 // CUDEVICEGETATTRIBUTE-NEXT:  /* 16 */ *pi = dpct::get_device(d).is_native_atomic_supported();
 
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuDeviceGetP2PAttribute | FileCheck %s -check-prefix=CUDEVICEGETP2PATTRIBUTE
+// CUDEVICEGETP2PATTRIBUTE: CUDA API:
+// CUDEVICEGETP2PATTRIBUTE-NEXT:   CUdevice_P2PAttribute attrib;
+// CUDEVICEGETP2PATTRIBUTE-NEXT:   /* 1 */ cuDeviceGetP2PAttribute(
+// CUDEVICEGETP2PATTRIBUTE-NEXT:       &value /*int **/, CU_DEVICE_P2P_ATTRIBUTE_ACCESS_SUPPORTED,
+// CUDEVICEGETP2PATTRIBUTE-NEXT:       srcDevice /*CUdevice*/, dstDevice /*CUdevice*/);
+// CUDEVICEGETP2PATTRIBUTE-NEXT:   /* 2 */ cuDeviceGetP2PAttribute(
+// CUDEVICEGETP2PATTRIBUTE-NEXT:       &value /*int **/, CU_DEVICE_P2P_ATTRIBUTE_NATIVE_ATOMIC_SUPPORTED,
+// CUDEVICEGETP2PATTRIBUTE-NEXT:       srcDevice /*CUdevice*/, dstDevice /*CUdevice*/);
+// CUDEVICEGETP2PATTRIBUTE-NEXT: Is migrated to:
+// CUDEVICEGETP2PATTRIBUTE-NEXT:   sycl::ext::oneapi::peer_access attrib;
+// CUDEVICEGETP2PATTRIBUTE-NEXT:   /* 1 */ value = dpct::get_device(srcDevice).ext_oneapi_can_access_peer(dpct::get_device(dstDevice), sycl::ext::oneapi::peer_access::access_supported);
+// CUDEVICEGETP2PATTRIBUTE-NEXT:   /* 2 */ value = dpct::get_device(srcDevice).ext_oneapi_can_access_peer(dpct::get_device(dstDevice), sycl::ext::oneapi::peer_access::atomics_supported);
+
 // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuDeviceGetCount | FileCheck %s -check-prefix=CUDEVICEGETCOUNT
 // CUDEVICEGETCOUNT: CUDA API:
 // CUDEVICEGETCOUNT-NEXT:   cuDeviceGetCount(pi /*int **/);
@@ -125,6 +139,19 @@
 // CUDEVICEPRIMARYCTXRELEASE-NEXT: The API is Removed.
 // CUDEVICEPRIMARYCTXRELEASE-EMPTY:
 
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuDevicePrimaryCtxSetFlags | FileCheck %s -check-prefix=CUDEVICEPRIMARYCTXSETFLAGS
+// CUDEVICEPRIMARYCTXSETFLAGS: CUDA API:
+// CUDEVICEPRIMARYCTXSETFLAGS-NEXT:   cuDevicePrimaryCtxSetFlags(device /*CUdevice*/, flags /*unsigned int*/);
+// CUDEVICEPRIMARYCTXSETFLAGS-NEXT: The API is Removed.
+// CUDEVICEPRIMARYCTXSETFLAGS-EMPTY:
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuDevicePrimaryCtxGetState | FileCheck %s -check-prefix=CUDEVICEPRIMARYCTXGETSTATE
+// CUDEVICEPRIMARYCTXGETSTATE: CUDA API:
+// CUDEVICEPRIMARYCTXGETSTATE-NEXT:   cuDevicePrimaryCtxGetState(device /*CUdevice*/, &flags /*unsigned int **/,
+// CUDEVICEPRIMARYCTXGETSTATE-NEXT:                              &active /*int **/);
+// CUDEVICEPRIMARYCTXGETSTATE-NEXT: The API is Removed.
+// CUDEVICEPRIMARYCTXGETSTATE-EMPTY:
+
 // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuDevicePrimaryCtxRetain | FileCheck %s -check-prefix=CUDEVICEPRIMARYCTXRETAIN
 // CUDEVICEPRIMARYCTXRETAIN: CUDA API:
 // CUDEVICEPRIMARYCTXRETAIN-NEXT:   cuDevicePrimaryCtxRetain(pc /*CUcontext **/, d /*CUdevice*/);

clang/test/dpct/query_api_mapping/Driver/test_memory_management.cu

@@ -106,6 +106,12 @@
 // CUMEMCPY2D-NEXT:   const dpct::memcpy_parameter *c;
 // CUMEMCPY2D-NEXT:   dpct::dpct_memcpy(*c);
 
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemcpy2DUnaligned | FileCheck %s -check-prefix=CUMEMCPY2DUNALIGNED
+// CUMEMCPY2DUNALIGNED: CUDA API:
+// CUMEMCPY2DUNALIGNED-NEXT:   cuMemcpy2DUnaligned(copy2d_desc /*const CUDA_MEMCPY2D **/);
+// CUMEMCPY2DUNALIGNED-NEXT: Is migrated to:
+// CUMEMCPY2DUNALIGNED-NEXT:   dpct::dpct_memcpy(*copy2d_desc /*const CUDA_MEMCPY2D **/);
+
 // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemcpy2DAsync | FileCheck %s -check-prefix=CUMEMCPY2DASYNC
 // CUMEMCPY2DASYNC: CUDA API:
 // CUMEMCPY2DASYNC-NEXT:   const CUDA_MEMCPY2D *c;

clang/test/dpct/query_api_mapping/nvSHMEM/test_library_setup_exit_query.cu

@@ -0,0 +1,56 @@
+// REQUIRES: system-linux
+// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=nvshmem_finalize | FileCheck %s -check-prefix=NVSHMEM_FINALIZE
+// NVSHMEM_FINALIZE: CUDA API:
+// NVSHMEM_FINALIZE-NEXT:   nvshmem_finalize(/*void*/);
+// NVSHMEM_FINALIZE-NEXT: Is migrated to:
+// NVSHMEM_FINALIZE-NEXT:   ishmem_finalize();
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=nvshmem_info_get_name | FileCheck %s -check-prefix=NVSHMEM_INFO_GET_NAME
+// NVSHMEM_INFO_GET_NAME: CUDA API:
+// NVSHMEM_INFO_GET_NAME-NEXT:   nvshmem_info_get_name(name /*char **/);
+// NVSHMEM_INFO_GET_NAME-NEXT: Is migrated to:
+// NVSHMEM_INFO_GET_NAME-NEXT:   ishmem_info_get_name(name);
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=nvshmem_info_get_version | FileCheck %s -check-prefix=NVSHMEM_INFO_GET_VERSION
+// NVSHMEM_INFO_GET_VERSION: CUDA API:
+// NVSHMEM_INFO_GET_VERSION-NEXT:   nvshmem_info_get_version(major /*int **/, minor /*int **/);
+// NVSHMEM_INFO_GET_VERSION-NEXT: Is migrated to:
+// NVSHMEM_INFO_GET_VERSION-NEXT:   ishmem_info_get_version(major, minor);
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=nvshmem_init | FileCheck %s -check-prefix=NVSHMEM_INIT
+// NVSHMEM_INIT: CUDA API:
+// NVSHMEM_INIT-NEXT:   nvshmem_init(/*void*/);
+// NVSHMEM_INIT-NEXT: Is migrated to:
+// NVSHMEM_INIT-NEXT:   ishmem_init();
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=nvshmem_my_pe | FileCheck %s -check-prefix=NVSHMEM_MY_PE
+// NVSHMEM_MY_PE: CUDA API:
+// NVSHMEM_MY_PE-NEXT:   nvshmem_my_pe(/*void*/);
+// NVSHMEM_MY_PE-NEXT: Is migrated to:
+// NVSHMEM_MY_PE-NEXT:   ishmem_my_pe();
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=nvshmem_n_pes | FileCheck %s -check-prefix=NVSHMEM_N_PES
+// NVSHMEM_N_PES: CUDA API:
+// NVSHMEM_N_PES-NEXT:   nvshmem_n_pes(/*void*/);
+// NVSHMEM_N_PES-NEXT: Is migrated to:
+// NVSHMEM_N_PES-NEXT:   ishmem_n_pes();
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=nvshmem_ptr | FileCheck %s -check-prefix=NVSHMEM_PTR
+// NVSHMEM_PTR: CUDA API:
+// NVSHMEM_PTR-NEXT:   local_ptr = nvshmem_ptr(dest /*const void **/, pe /*int*/);
+// NVSHMEM_PTR-NEXT: Is migrated to:
+// NVSHMEM_PTR-NEXT:   local_ptr = ishmem_ptr(dest, pe);
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=nvshmemx_init_attr | FileCheck %s -check-prefix=NVSHMEMX_INIT_ATTR
+// NVSHMEMX_INIT_ATTR: CUDA API:
+// NVSHMEMX_INIT_ATTR-NEXT:   nvshmemx_init_attr_t attributes;
+// NVSHMEMX_INIT_ATTR-NEXT:   /* 1 */ nvshmemx_init_attr(NVSHMEMX_INIT_WITH_MPI_COMM,
+// NVSHMEMX_INIT_ATTR-NEXT:                              &attributes /*nvshmemx_init_attr_t **/);
+// NVSHMEMX_INIT_ATTR-NEXT:   /* 2 */ nvshmemx_init_attr(NVSHMEMX_INIT_WITH_SHMEM,
+// NVSHMEMX_INIT_ATTR-NEXT:                              &attributes /*nvshmemx_init_attr_t **/);
+// NVSHMEMX_INIT_ATTR-NEXT: Is migrated to:
+// NVSHMEMX_INIT_ATTR-NEXT:   ishmemx_attr_t attributes;
+// NVSHMEMX_INIT_ATTR-NEXT:   /* 1 */ dpct::shmemx::init_attr(dpct::shmemx::RUNTIME_MPI, &attributes);
+// NVSHMEMX_INIT_ATTR-NEXT:   /* 2 */ dpct::shmemx::init_attr(dpct::shmemx::RUNTIME_OPENSHMEM, &attributes);

clang/test/dpct/query_api_mapping/nvSHMEM/test_mem_mgmt.cu

@@ -0,0 +1,26 @@
+// REQUIRES: system-linux
+// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=nvshmem_align | FileCheck %s -check-prefix=NVSHMEM_ALIGN
+// NVSHMEM_ALIGN: CUDA API:
+// NVSHMEM_ALIGN-NEXT:   nvshmem_align(alignment /*size_t*/, size /*size_t*/);
+// NVSHMEM_ALIGN-NEXT: Is migrated to:
+// NVSHMEM_ALIGN-NEXT:   ishmem_align(alignment, size);
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=nvshmem_calloc | FileCheck %s -check-prefix=NVSHMEM_CALLOC
+// NVSHMEM_CALLOC: CUDA API:
+// NVSHMEM_CALLOC-NEXT:   nvshmem_calloc(count /*size_t*/, size /*size_t*/);
+// NVSHMEM_CALLOC-NEXT: Is migrated to:
+// NVSHMEM_CALLOC-NEXT:   ishmem_calloc(count, size);
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=nvshmem_free | FileCheck %s -check-prefix=NVSHMEM_FREE
+// NVSHMEM_FREE: CUDA API:
+// NVSHMEM_FREE-NEXT:   nvshmem_free(ptr /*void **/);
+// NVSHMEM_FREE-NEXT: Is migrated to:
+// NVSHMEM_FREE-NEXT:   ishmem_free(ptr);
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=nvshmem_malloc | FileCheck %s -check-prefix=NVSHMEM_MALLOC
+// NVSHMEM_MALLOC: CUDA API:
+// NVSHMEM_MALLOC-NEXT:   nvshmem_malloc(size /*size_t*/);
+// NVSHMEM_MALLOC-NEXT: Is migrated to:
+// NVSHMEM_MALLOC-NEXT:   ishmem_malloc(size);