diff --git a/clang/examples/DPCT/Driver/cuArray3DGetDescriptor.cu b/clang/examples/DPCT/Driver/cuArray3DGetDescriptor.cu new file mode 100644 index 000000000000..4bd3053c7d84 --- /dev/null +++ b/clang/examples/DPCT/Driver/cuArray3DGetDescriptor.cu @@ -0,0 +1,7 @@ +// Option: --use-experimental-features=bindless_images + +void test(CUDA_ARRAY3D_DESCRIPTOR *desc, CUarray array) { + // Start + cuArray3DGetDescriptor(desc /*CUDA_ARRAY3D_DESCRIPTOR **/, array /*CUarray*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuArrayGetDescriptor.cu b/clang/examples/DPCT/Driver/cuArrayGetDescriptor.cu new file mode 100644 index 000000000000..8c554fc188c5 --- /dev/null +++ b/clang/examples/DPCT/Driver/cuArrayGetDescriptor.cu @@ -0,0 +1,7 @@ +// Option: --use-experimental-features=bindless_images + +void test(CUDA_ARRAY_DESCRIPTOR *desc, CUarray array) { + // Start + cuArrayGetDescriptor(desc /*CUDA_ARRAY_DESCRIPTOR **/, array /*CUarray*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuCtxCreate_v3.cu b/clang/examples/DPCT/Driver/cuCtxCreate_v3.cu new file mode 100644 index 000000000000..14bc03f7d6bf --- /dev/null +++ b/clang/examples/DPCT/Driver/cuCtxCreate_v3.cu @@ -0,0 +1,6 @@ +void test(CUcontext *ctx, CUexecAffinityParam *params_array, int num, + unsigned int flags, CUdevice device) { + // Start + cuCtxCreate_v3(ctx /*CUcontext **/, params_array /*CUexecAffinityParam **/, num /*int*/, flags /*unsigned int*/, device /*CUdevice*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuCtxCreate_v4.cu b/clang/examples/DPCT/Driver/cuCtxCreate_v4.cu new file mode 100644 index 000000000000..e2e7fc4fc337 --- /dev/null +++ b/clang/examples/DPCT/Driver/cuCtxCreate_v4.cu @@ -0,0 +1,6 @@ +void test(CUcontext *ctx, CUctxCreateParams *params_array, unsigned int flags, + CUdevice device) { + // Start + cuCtxCreate_v4(ctx /*CUcontext **/, params_array /*CUctxCreateParams **/, flags /*unsigned int*/, device /*CUdevice*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuGetErrorName.cu b/clang/examples/DPCT/Driver/cuGetErrorName.cu new file mode 100644 index 000000000000..472d44b78f8d --- /dev/null +++ b/clang/examples/DPCT/Driver/cuGetErrorName.cu @@ -0,0 +1,5 @@ +void test(CUresult r, const char **pstr) { + // Start + cuGetErrorName(r /*CUresult*/, pstr /*const char ***/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuMipmappedArrayCreate.cu b/clang/examples/DPCT/Driver/cuMipmappedArrayCreate.cu new file mode 100644 index 000000000000..8e51e5c919d8 --- /dev/null +++ b/clang/examples/DPCT/Driver/cuMipmappedArrayCreate.cu @@ -0,0 +1,8 @@ +// Option: --use-experimental-features=bindless_images + +void test(CUmipmappedArray *array, CUDA_ARRAY3D_DESCRIPTOR *desc, + unsigned int levels) { + // Start + cuMipmappedArrayCreate(array /*CUmipmappedArray **/, desc /*CUDA_ARRAY3D_DESCRIPTOR **/, levels /*unsigned int*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Guide/surf2DLayeredwrite.cu b/clang/examples/DPCT/Guide/surf2DLayeredwrite.cu new file mode 100644 index 000000000000..4e6853a1d15b --- /dev/null +++ b/clang/examples/DPCT/Guide/surf2DLayeredwrite.cu @@ -0,0 +1,9 @@ +// Option: --use-experimental-features=bindless_images +template +__global__ void test(float data, cudaSurfaceObject_t surf, int x, int y, + int layer) { + // Start + surf2DLayeredwrite(data /*float*/, surf /*cudaSurfaceObject_t*/, x /*int*/, + y /*int*/, layer /*int*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Guide/surf2Dread.cu b/clang/examples/DPCT/Guide/surf2Dread.cu new file mode 100644 index 000000000000..a7922f2f31aa --- /dev/null +++ b/clang/examples/DPCT/Guide/surf2Dread.cu @@ -0,0 +1,8 @@ +// Option: --use-experimental-features=bindless_images +template +__global__ void test(float data, cudaSurfaceObject_t surf, int x, int y) { + // Start + data /*float*/ = + surf2Dread(surf /*cudaSurfaceObject_t*/, x /*int*/, y /*int*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Guide/surf2Dwrite.cu b/clang/examples/DPCT/Guide/surf2Dwrite.cu new file mode 100644 index 000000000000..91c06b10d0cc --- /dev/null +++ b/clang/examples/DPCT/Guide/surf2Dwrite.cu @@ -0,0 +1,8 @@ +// Option: --use-experimental-features=bindless_images +template +__global__ void test(float data, cudaSurfaceObject_t surf, int x, int y) { + // Start + surf2Dwrite(data /*float*/, surf /*cudaSurfaceObject_t*/, x /*int*/, + y /*int*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Guide/surf3Dread.cu b/clang/examples/DPCT/Guide/surf3Dread.cu new file mode 100644 index 000000000000..ecd106e7ea61 --- /dev/null +++ b/clang/examples/DPCT/Guide/surf3Dread.cu @@ -0,0 +1,9 @@ +// Option: --use-experimental-features=bindless_images +template +__global__ void test(float data, cudaSurfaceObject_t surf, int x, int y, + int z) { + // Start + data /*float*/ = surf3Dread(surf /*cudaSurfaceObject_t*/, x /*int*/, + y /*int*/, z /*int*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Guide/surf3Dwrite.cu b/clang/examples/DPCT/Guide/surf3Dwrite.cu new file mode 100644 index 000000000000..17f8b6d552dd --- /dev/null +++ b/clang/examples/DPCT/Guide/surf3Dwrite.cu @@ -0,0 +1,9 @@ +// Option: --use-experimental-features=bindless_images +template +__global__ void test(float data, cudaSurfaceObject_t surf, int x, int y, + int z) { + // Start + surf3Dwrite(data /*float*/, surf /*cudaSurfaceObject_t*/, x /*int*/, + y /*int*/, z /*int*/); + // End +} \ No newline at end of file diff --git a/clang/test/dpct/query_api_mapping/Driver/test-after12.4.cu b/clang/test/dpct/query_api_mapping/Driver/test-after12.4.cu new file mode 100644 index 000000000000..1e2bf0c9f1ef --- /dev/null +++ b/clang/test/dpct/query_api_mapping/Driver/test-after12.4.cu @@ -0,0 +1,14 @@ +// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2, 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 +// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2, 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 + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuCtxCreate_v3 | FileCheck %s -check-prefix=CUCTXCREATE_V3 +// CUCTXCREATE_V3: CUDA API: +// CUCTXCREATE_V3-NEXT: cuCtxCreate_v3(ctx /*CUcontext **/, params_array /*CUexecAffinityParam **/, num /*int*/, flags /*unsigned int*/, device /*CUdevice*/); +// CUCTXCREATE_V3-NEXT: Is migrated to: +// CUCTXCREATE_V3-NEXT: *ctx = dpct::push_device_for_curr_thread(device); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuCtxCreate_v4 | FileCheck %s -check-prefix=CUCTXCREATE_V4 +// CUCTXCREATE_V4: CUDA API: +// CUCTXCREATE_V4-NEXT: cuCtxCreate_v4(ctx /*CUcontext **/, params_array /*CUctxCreateParams **/, flags /*unsigned int*/, device /*CUdevice*/); +// CUCTXCREATE_V4-NEXT: Is migrated to: +// CUCTXCREATE_V4-NEXT: *ctx = dpct::push_device_for_curr_thread(device); diff --git a/clang/test/dpct/query_api_mapping/Driver/test.cu b/clang/test/dpct/query_api_mapping/Driver/test.cu index f62b789db3fb..2455ef9b3fcf 100644 --- a/clang/test/dpct/query_api_mapping/Driver/test.cu +++ b/clang/test/dpct/query_api_mapping/Driver/test.cu @@ -9,6 +9,15 @@ // CUGETERRORSTRING-NEXT: */ // CUGETERRORSTRING-NEXT: *ppc = dpct::get_error_string_dummy(r); +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuGetErrorName | FileCheck %s -check-prefix=CUGETERRORNAME +// CUGETERRORNAME: CUDA API: +// CUGETERRORNAME-NEXT: cuGetErrorName(r /*CUresult*/, pstr /*const char ***/); +// CUGETERRORNAME-NEXT: Is migrated to: +// CUGETERRORNAME-NEXT: /* +// CUGETERRORNAME-NEXT: DPCT1009:0: SYCL reports errors using exceptions and does not use error codes. Please replace the "get_error_string_dummy(...)" with a real error-handling function. +// CUGETERRORNAME-NEXT: */ +// CUGETERRORNAME-NEXT: *pstr = dpct::get_error_string_dummy(r); + /// Initialization // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuInit | FileCheck %s -check-prefix=CUINIT @@ -441,6 +450,24 @@ // CUTEXREFGETADDRESSMODE-NEXT: dpct::image_wrapper_base_p t; // CUTEXREFGETADDRESSMODE-NEXT: *pa = t->get_addressing_mode(); +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuArray3DGetDescriptor | FileCheck %s -check-prefix=CUARRAY3DGETDESCRIPTOR +// CUARRAY3DGETDESCRIPTOR: CUDA API: +// CUARRAY3DGETDESCRIPTOR-NEXT: cuArray3DGetDescriptor(desc /*CUDA_ARRAY3D_DESCRIPTOR **/, array /*CUarray*/); +// CUARRAY3DGETDESCRIPTOR-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// CUARRAY3DGETDESCRIPTOR-NEXT: *desc = array->get_desc(); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuArrayGetDescriptor | FileCheck %s -check-prefix=CUARRAYGETDESCRIPTOR +// CUARRAYGETDESCRIPTOR: CUDA API: +// CUARRAYGETDESCRIPTOR-NEXT: cuArrayGetDescriptor(desc /*CUDA_ARRAY_DESCRIPTOR **/, array /*CUarray*/); +// CUARRAYGETDESCRIPTOR-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// CUARRAYGETDESCRIPTOR-NEXT: *desc = array->get_desc(); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMipmappedArrayCreate | FileCheck %s -check-prefix=CUMIPMAPPEDARRAYCREATE +// CUMIPMAPPEDARRAYCREATE: CUDA API: +// CUMIPMAPPEDARRAYCREATE-NEXT: cuMipmappedArrayCreate(array /*CUmipmappedArray **/, desc /*CUDA_ARRAY3D_DESCRIPTOR **/, levels /*unsigned int*/); +// CUMIPMAPPEDARRAYCREATE-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// CUMIPMAPPEDARRAYCREATE-NEXT: *array = new dpct::experimental::image_mem_wrapper(*desc, levels); + // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuTexRefGetFilterMode | FileCheck %s -check-prefix=CUTEXREFGETFILTERMODE // CUTEXREFGETFILTERMODE: CUDA API: // CUTEXREFGETFILTERMODE-NEXT: CUtexref t; diff --git a/clang/test/dpct/query_api_mapping/Guide/test.cu b/clang/test/dpct/query_api_mapping/Guide/test.cu index ed297eab73f7..3c377537b24d 100644 --- a/clang/test/dpct/query_api_mapping/Guide/test.cu +++ b/clang/test/dpct/query_api_mapping/Guide/test.cu @@ -80,6 +80,42 @@ // TEX2DLOD-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): // TEX2DLOD-NEXT: sycl::ext::oneapi::experimental::sample_mipmap(t, sycl::float2(f1, f2), f3); +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=surf2DLayeredwrite | FileCheck %s -check-prefix=SURF2DLAYEREDWRITE +// SURF2DLAYEREDWRITE: CUDA API: +// SURF2DLAYEREDWRITE-NEXT: surf2DLayeredwrite(data /*float*/, surf /*cudaSurfaceObject_t*/, x /*int*/, +// SURF2DLAYEREDWRITE-NEXT: y /*int*/, layer /*int*/); +// SURF2DLAYEREDWRITE-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// SURF2DLAYEREDWRITE-NEXT: sycl::ext::oneapi::experimental::write_image_array(surf, sycl::int2(x / sizeof(data), y), layer, data); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=surf2Dread | FileCheck %s -check-prefix=SURF2DREAD +// SURF2DREAD: CUDA API: +// SURF2DREAD-NEXT: data /*float*/ = +// SURF2DREAD-NEXT: surf2Dread(surf /*cudaSurfaceObject_t*/, x /*int*/, y /*int*/); +// SURF2DREAD-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// SURF2DREAD-NEXT: data /*float*/ = +// SURF2DREAD-NEXT: dpct::experimental::fetch_image_by_byte(surf, sycl::int2(x, y)); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=surf2Dwrite | FileCheck %s -check-prefix=SURF2DWRITE +// SURF2DWRITE: CUDA API: +// SURF2DWRITE-NEXT: surf2Dwrite(data /*float*/, surf /*cudaSurfaceObject_t*/, x /*int*/, +// SURF2DWRITE-NEXT: y /*int*/); +// SURF2DWRITE-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// SURF2DWRITE-NEXT: sycl::ext::oneapi::experimental::write_image(surf, sycl::int2(x / sizeof(data), y), data); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=surf3Dread | FileCheck %s -check-prefix=SURF3DREAD +// SURF3DREAD: CUDA API: +// SURF3DREAD-NEXT: data /*float*/ = surf3Dread(surf /*cudaSurfaceObject_t*/, x /*int*/, +// SURF3DREAD-NEXT: y /*int*/, z /*int*/); +// SURF3DREAD-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// SURF3DREAD-NEXT: data /*float*/ = dpct::experimental::fetch_image_by_byte(surf, sycl::int3(x, y, z)); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=surf3Dwrite | FileCheck %s -check-prefix=SURF3DWRITE +// SURF3DWRITE: CUDA API: +// SURF3DWRITE-NEXT: surf3Dwrite(data /*float*/, surf /*cudaSurfaceObject_t*/, x /*int*/, +// SURF3DWRITE-NEXT: y /*int*/, z /*int*/); +// SURF3DWRITE-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// SURF3DWRITE-NEXT: sycl::ext::oneapi::experimental::write_image(surf, sycl::int3(x / sizeof(data), y, z), data); + // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=tex3D | FileCheck %s -check-prefix=TEX3D // TEX3D: CUDA API: // TEX3D-NEXT: tex3D(t /*cudaTextureObject_t*/, f1 /*float*/, f2 /*float*/, f3 /*float*/); diff --git a/clang/test/dpct/query_api_mapping/test_all.cu b/clang/test/dpct/query_api_mapping/test_all.cu index 0b0d5dc9c58d..282d160017d5 100644 --- a/clang/test/dpct/query_api_mapping/test_all.cu +++ b/clang/test/dpct/query_api_mapping/test_all.cu @@ -679,8 +679,10 @@ // CHECK-NEXT: cospi // CHECK-NEXT: cospif // CHECK-NEXT: cuArray3DCreate +// CHECK-NEXT: cuArray3DGetDescriptor // CHECK-NEXT: cuArrayCreate // CHECK-NEXT: cuArrayDestroy +// CHECK-NEXT: cuArrayGetDescriptor // CHECK-NEXT: cuCabs // CHECK-NEXT: cuCabsf // CHECK-NEXT: cuCadd @@ -702,6 +704,8 @@ // CHECK-NEXT: cuCsub // CHECK-NEXT: cuCsubf // CHECK-NEXT: cuCtxCreate +// CHECK-NEXT: cuCtxCreate_v3 +// CHECK-NEXT: cuCtxCreate_v4 // CHECK-NEXT: cuCtxDestroy // CHECK-NEXT: cuCtxEnablePeerAccess // CHECK-NEXT: cuCtxGetApiVersion @@ -739,6 +743,7 @@ // CHECK-NEXT: cuFuncGetAttribute // CHECK-NEXT: cuFuncSetAttribute // CHECK-NEXT: cuFuncSetCacheConfig +// CHECK-NEXT: cuGetErrorName // CHECK-NEXT: cuGetErrorString // CHECK-NEXT: cuGraphicsMapResources // CHECK-NEXT: cuGraphicsResourceGetMappedPointer @@ -805,6 +810,7 @@ // CHECK-NEXT: cuMemsetD32Async // CHECK-NEXT: cuMemsetD8 // CHECK-NEXT: cuMemsetD8Async +// CHECK-NEXT: cuMipmappedArrayCreate // CHECK-NEXT: cuMipmappedArrayDestroy // CHECK-NEXT: cuMipmappedArrayGetLevel // CHECK-NEXT: cuModuleGetFunction @@ -2442,6 +2448,11 @@ // CHECK-NEXT: sqrtf // CHECK-NEXT: surf1Dread // CHECK-NEXT: surf1Dwrite +// CHECK-NEXT: surf2DLayeredwrite +// CHECK-NEXT: surf2Dread +// CHECK-NEXT: surf2Dwrite +// CHECK-NEXT: surf3Dread +// CHECK-NEXT: surf3Dwrite // CHECK-NEXT: tan // CHECK-NEXT: tanf // CHECK-NEXT: tanh