diff --git a/clang/examples/DPCT/Runtime/cudaGraphExecDestroy.cu b/clang/examples/DPCT/Runtime/cudaGraphExecDestroy.cu new file mode 100644 index 000000000000..9886befde6f2 --- /dev/null +++ b/clang/examples/DPCT/Runtime/cudaGraphExecDestroy.cu @@ -0,0 +1,7 @@ +// Option: --use-experimental-features=graph + +void test(cudaGraphExec_t graph_exec) { + // Start + cudaGraphExecDestroy(graph_exec /*cudaGraphExec_t*/); + // End +} diff --git a/clang/examples/DPCT/Runtime/cudaGraphExecUpdate.cu b/clang/examples/DPCT/Runtime/cudaGraphExecUpdate.cu new file mode 100644 index 000000000000..58e41c8fdcf4 --- /dev/null +++ b/clang/examples/DPCT/Runtime/cudaGraphExecUpdate.cu @@ -0,0 +1,10 @@ +// Option: --use-experimental-features=graph + +void test(cudaGraphExec_t graph_exec, cudaGraph_t graph, cudaGraphNode_t *node, + cudaGraphExecUpdateResult *result) { + // Start + cudaGraphExecUpdate(graph_exec /*cudaGraphExec_t*/, graph /*cudaGraph_t*/, + node /*cudaGraphNode_t **/, + result /*cudaGraphExecUpdateResult **/); + // End +} diff --git a/clang/examples/DPCT/Runtime/cudaGraphGetNodes.cu b/clang/examples/DPCT/Runtime/cudaGraphGetNodes.cu new file mode 100644 index 000000000000..2ad4bd05d1ad --- /dev/null +++ b/clang/examples/DPCT/Runtime/cudaGraphGetNodes.cu @@ -0,0 +1,8 @@ +// Option: --use-experimental-features=graph + +void test(cudaGraph_t graph, cudaGraphNode_t *nodes, size_t *num_nodes) { + // Start + cudaGraphGetNodes(graph /*cudaGraph_t*/, nodes /*cudaGraphNode_t **/, + num_nodes /*size_t **/); + // End +} diff --git a/clang/examples/DPCT/Runtime/cudaGraphGetRootNodes.cu b/clang/examples/DPCT/Runtime/cudaGraphGetRootNodes.cu new file mode 100644 index 000000000000..c6b1f9349029 --- /dev/null +++ b/clang/examples/DPCT/Runtime/cudaGraphGetRootNodes.cu @@ -0,0 +1,8 @@ +// Option: --use-experimental-features=graph + +void test(cudaGraph_t graph, cudaGraphNode_t *nodes, size_t *num_nodes) { + // Start + cudaGraphGetRootNodes(graph /*cudaGraph_t*/, nodes /*cudaGraphNode_t **/, + num_nodes /*size_t **/); + // End +} diff --git a/clang/examples/DPCT/Runtime/cudaGraphInstantiate.cu b/clang/examples/DPCT/Runtime/cudaGraphInstantiate.cu new file mode 100644 index 000000000000..778c59b2f22e --- /dev/null +++ b/clang/examples/DPCT/Runtime/cudaGraphInstantiate.cu @@ -0,0 +1,10 @@ +// Option: --use-experimental-features=graph + +void test(cudaGraphExec_t *graph_exec, cudaGraph_t graph, cudaGraphNode_t *node, + char *buffer, size_t buffer_size) { + // Start + cudaGraphInstantiate(graph_exec /*cudaGraphExec_t **/, graph /*cudaGraph_t*/, + node /*cudaGraphNode_t **/, buffer /*char **/, + buffer_size /*size_t*/); + // End +} diff --git a/clang/examples/DPCT/Runtime/cudaGraphLaunch.cu b/clang/examples/DPCT/Runtime/cudaGraphLaunch.cu new file mode 100644 index 000000000000..3c5c7ce328d5 --- /dev/null +++ b/clang/examples/DPCT/Runtime/cudaGraphLaunch.cu @@ -0,0 +1,7 @@ +// Option: --use-experimental-features=graph + +void test(cudaGraphExec_t graph_exec, cudaStream_t stream) { + // Start + cudaGraphLaunch(graph_exec /*cudaGraphExec_t*/, stream /*cudaStream_t*/); + // End +} diff --git a/clang/examples/DPCT/Runtime/cudaGraphNodeGetType.cu b/clang/examples/DPCT/Runtime/cudaGraphNodeGetType.cu new file mode 100644 index 000000000000..f081af0e9ba3 --- /dev/null +++ b/clang/examples/DPCT/Runtime/cudaGraphNodeGetType.cu @@ -0,0 +1,8 @@ +// Option: --use-experimental-features=graph + +void test(cudaGraphNode_t node, cudaGraphNodeType *node_type) { + // Start + cudaGraphNodeGetType(node /*cudaGraphNode_t*/, + node_type /*cudaGraphNodeType **/); + // End +} diff --git a/clang/examples/DPCT/Runtime/surf1Dread.cu b/clang/examples/DPCT/Runtime/surf1Dread.cu new file mode 100644 index 000000000000..9bd423cabd61 --- /dev/null +++ b/clang/examples/DPCT/Runtime/surf1Dread.cu @@ -0,0 +1,12 @@ +// Option: --use-experimental-features=bindless_images + +template +__global__ void test(T *ptr, int x, cudaSurfaceBoundaryMode boundary_mode, + cudaSurfaceObject_t obj) { + // Start + surf1Dread(obj /*cudaSurfaceObject_t*/, x /*int*/, + boundary_mode /*cudaSurfaceBoundaryMode*/); + surf1Dread(ptr /*T **/, obj /*cudaSurfaceObject_t*/, x /*int*/, + boundary_mode /*cudaSurfaceBoundaryMode*/); + // End +} diff --git a/clang/examples/DPCT/Runtime/surf1Dwrite.cu b/clang/examples/DPCT/Runtime/surf1Dwrite.cu new file mode 100644 index 000000000000..73c021271b91 --- /dev/null +++ b/clang/examples/DPCT/Runtime/surf1Dwrite.cu @@ -0,0 +1,10 @@ +// Option: --use-experimental-features=bindless_images + +template +__global__ void test(T val, int x, cudaSurfaceBoundaryMode boundary_mode, + cudaSurfaceObject_t obj) { + // Start + surf1Dwrite(val /*T*/, obj /*cudaSurfaceObject_t*/, x /*int*/, + boundary_mode /*cudaSurfaceBoundaryMode*/); + // End +} diff --git a/clang/examples/DPCT/Runtime/tex1DLayered.cu b/clang/examples/DPCT/Runtime/tex1DLayered.cu new file mode 100644 index 000000000000..606d5ec55b44 --- /dev/null +++ b/clang/examples/DPCT/Runtime/tex1DLayered.cu @@ -0,0 +1,10 @@ +// Option: --use-experimental-features=bindless_images + +template +__global__ void test(T *ptr, float x, int layer, cudaTextureObject_t tex) { + // Start + tex1DLayered(tex /*cudaTextureObject_t*/, x /*float*/, layer /*int*/); + tex1DLayered(ptr /*T **/, tex /*cudaTextureObject_t*/, x /*float*/, + layer /*int*/); + // End +} diff --git a/clang/examples/DPCT/Runtime/tex2DLayered.cu b/clang/examples/DPCT/Runtime/tex2DLayered.cu new file mode 100644 index 000000000000..74bfe0c96ef2 --- /dev/null +++ b/clang/examples/DPCT/Runtime/tex2DLayered.cu @@ -0,0 +1,16 @@ +// Option: --use-experimental-features=bindless_images + +template +__global__ void test(T *ptr, float x, float y, int layer, + cudaTextureObject_t tex, bool *is_resident) { + // Start + tex2DLayered(tex /*cudaTextureObject_t*/, x /*float*/, y /*float*/, + layer /*int*/); + tex2DLayered(ptr /*T **/, tex /*cudaTextureObject_t*/, x /*float*/, + y /*float*/, layer /*int*/); + tex2DLayered(tex /*cudaTextureObject_t*/, x /*float*/, y /*float*/, + layer /*int*/, is_resident /*bool **/); + tex2DLayered(ptr /*T **/, tex /*cudaTextureObject_t*/, x /*float*/, + y /*float*/, layer /*int*/, is_resident /*bool **/); + // End +} diff --git a/clang/test/dpct/query_api_mapping/Runtime/test-after10.cu b/clang/test/dpct/query_api_mapping/Runtime/test-after10.cu index 115226d39b50..517773b198fc 100644 --- a/clang/test/dpct/query_api_mapping/Runtime/test-after10.cu +++ b/clang/test/dpct/query_api_mapping/Runtime/test-after10.cu @@ -1,14 +1,6 @@ // UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2 // UNSUPPORTED: v8.0, v9.0, v9.1, v9.2 -/// Stream Management - -// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cudaStreamBeginCapture | FileCheck %s -check-prefix=CUDASTREAMBEGINCAPTURE -// CUDASTREAMBEGINCAPTURE: CUDA API: -// CUDASTREAMBEGINCAPTURE-NEXT: cudaStreamBeginCapture(s /*cudaStream_t*/, sc /*cudaStreamCaptureMode*/); -// CUDASTREAMBEGINCAPTURE-NEXT: Is migrated to (with the option --use-experimental-features=graph): -// CUDASTREAMBEGINCAPTURE-NEXT: dpct::experimental::begin_recording(s); - // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cudaStreamEndCapture | FileCheck %s -check-prefix=CUDASTREAMENDCAPTURE // CUDASTREAMENDCAPTURE: CUDA API: // CUDASTREAMENDCAPTURE-NEXT: cudaStreamEndCapture(s /*cudaStream_t*/, pg /*cudaGraph_t **/); @@ -22,8 +14,6 @@ // CUDASTREAMISCAPTURING-NEXT: Is migrated to (with the option --use-experimental-features=graph): // CUDASTREAMISCAPTURING-NEXT: *ps = s->ext_oneapi_get_state(); - - // RUN: dpct --cuda-include-path="%cuda-path/include" -query-api-mapping=cudaLaunchHostFunc | FileCheck %s -check-prefix=CUDALAUNCHHOSTFUNC // CUDALAUNCHHOSTFUNC: CUDA API: // CUDALAUNCHHOSTFUNC-NEXT: cudaLaunchHostFunc(stream/*cudaStream_t*/, fn/*cudaHostFn_t*/, userData/*void**/); @@ -32,4 +22,88 @@ // CUDALAUNCHHOSTFUNC-NEXT: cgh.host_task([=](){ // CUDALAUNCHHOSTFUNC-NEXT: fn(userData); // CUDALAUNCHHOSTFUNC-NEXT: }); -// CUDALAUNCHHOSTFUNC-NEXT: }); \ No newline at end of file +// CUDALAUNCHHOSTFUNC-NEXT: }); + + +// RUN: dpct --cuda-include-path="%cuda-path/include" -query-api-mapping=cudaGraphExecDestroy | FileCheck %s -check-prefix=cudaGraphExecDestroy +// cudaGraphExecDestroy: CUDA API: +// cudaGraphExecDestroy-NEXT: cudaGraphExecDestroy(graph_exec /*cudaGraphExec_t*/); +// cudaGraphExecDestroy-NEXT: Is migrated to (with the option --use-experimental-features=graph): +// cudaGraphExecDestroy-NEXT: delete (graph_exec); + +// RUN: dpct --cuda-include-path="%cuda-path/include" -query-api-mapping=cudaGraphGetNodes | FileCheck %s -check-prefix=cudaGraphGetNodes +// cudaGraphGetNodes: CUDA API: +// cudaGraphGetNodes-NEXT: cudaGraphGetNodes(graph /*cudaGraph_t*/, nodes /*cudaGraphNode_t **/, +// cudaGraphGetNodes-NEXT: num_nodes /*size_t **/); +// cudaGraphGetNodes-NEXT: Is migrated to (with the option --use-experimental-features=graph): +// cudaGraphGetNodes-NEXT: dpct::experimental::get_nodes(graph, nodes, num_nodes); + +// RUN: dpct --cuda-include-path="%cuda-path/include" -query-api-mapping=cudaGraphGetRootNodes | FileCheck %s -check-prefix=cudaGraphGetRootNodes +// cudaGraphGetRootNodes: CUDA API: +// cudaGraphGetRootNodes-NEXT: cudaGraphGetRootNodes(graph /*cudaGraph_t*/, nodes /*cudaGraphNode_t **/, +// cudaGraphGetRootNodes-NEXT: num_nodes /*size_t **/); +// cudaGraphGetRootNodes-NEXT: Is migrated to (with the option --use-experimental-features=graph): +// cudaGraphGetRootNodes-NEXT: dpct::experimental::get_root_nodes(graph, nodes, num_nodes); + +// RUN: dpct --cuda-include-path="%cuda-path/include" -query-api-mapping=cudaGraphInstantiate | FileCheck %s -check-prefix=cudaGraphInstantiate +// cudaGraphInstantiate: CUDA API: +// cudaGraphInstantiate-NEXT: cudaGraphInstantiate(graph_exec /*cudaGraphExec_t **/, graph /*cudaGraph_t*/, +// cudaGraphInstantiate-NEXT: node /*cudaGraphNode_t **/, buffer /*char **/, +// cudaGraphInstantiate-NEXT: buffer_size /*size_t*/); +// cudaGraphInstantiate-NEXT: Is migrated to (with the option --use-experimental-features=graph): +// cudaGraphInstantiate-NEXT: *graph_exec = new sycl::ext::oneapi::experimental::command_graph(graph->finalize()); + +// RUN: dpct --cuda-include-path="%cuda-path/include" -query-api-mapping=cudaGraphLaunch | FileCheck %s -check-prefix=cudaGraphLaunch +// cudaGraphLaunch: CUDA API: +// cudaGraphLaunch-NEXT: cudaGraphLaunch(graph_exec /*cudaGraphExec_t*/, stream /*cudaStream_t*/); +// cudaGraphLaunch-NEXT: Is migrated to (with the option --use-experimental-features=graph): +// cudaGraphLaunch-NEXT: stream->ext_oneapi_graph(*graph_exec); + +// RUN: dpct --cuda-include-path="%cuda-path/include" -query-api-mapping=cudaGraphNodeGetType | FileCheck %s -check-prefix=cudaGraphNodeGetType +// cudaGraphNodeGetType: CUDA API: +// cudaGraphNodeGetType-NEXT: cudaGraphNodeGetType(node /*cudaGraphNode_t*/, +// cudaGraphNodeGetType-NEXT: node_type /*cudaGraphNodeType **/); +// cudaGraphNodeGetType-NEXT: Is migrated to (with the option --use-experimental-features=graph): +// cudaGraphNodeGetType-NEXT: *node_type = node->get_type(); + +// RUN: dpct --cuda-include-path="%cuda-path/include" -query-api-mapping=surf1Dread | FileCheck %s -check-prefix=surf1Dread +// surf1Dread: CUDA API: +// surf1Dread-NEXT: surf1Dread(obj /*cudaSurfaceObject_t*/, x /*int*/, +// surf1Dread-NEXT: boundary_mode /*cudaSurfaceBoundaryMode*/); +// surf1Dread-NEXT: surf1Dread(ptr /*T **/, obj /*cudaSurfaceObject_t*/, x /*int*/, +// surf1Dread-NEXT: boundary_mode /*cudaSurfaceBoundaryMode*/); +// surf1Dread-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// surf1Dread-NEXT: dpct::experimental::fetch_image_by_byte(obj, int(x)); +// surf1Dread-NEXT: *ptr = dpct::experimental::fetch_image_by_byte(obj, int(x)); + +// RUN: dpct --cuda-include-path="%cuda-path/include" -query-api-mapping=surf1Dwrite | FileCheck %s -check-prefix=surf1Dwrite +// surf1Dwrite: CUDA API: +// surf1Dwrite-NEXT: surf1Dwrite(val /*T*/, obj /*cudaSurfaceObject_t*/, x /*int*/, +// surf1Dwrite-NEXT: boundary_mode /*cudaSurfaceBoundaryMode*/); +// surf1Dwrite-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// surf1Dwrite-NEXT: sycl::ext::oneapi::experimental::write_image(obj, int(x / sizeof(val)), val); + +// RUN: dpct --cuda-include-path="%cuda-path/include" -query-api-mapping=tex1DLayered | FileCheck %s -check-prefix=tex1DLayered +// tex1DLayered: CUDA API: +// tex1DLayered-NEXT: tex1DLayered(tex /*cudaTextureObject_t*/, x /*float*/, layer /*int*/); +// tex1DLayered-NEXT: tex1DLayered(ptr /*T **/, tex /*cudaTextureObject_t*/, x /*float*/, +// tex1DLayered-NEXT: layer /*int*/); +// tex1DLayered-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// tex1DLayered-NEXT: sycl::ext::oneapi::experimental::sample_image_array(tex, float(x), layer); +// tex1DLayered-NEXT: *ptr = sycl::ext::oneapi::experimental::sample_image_array(tex, float(x), layer); + +// RUN: dpct --cuda-include-path="%cuda-path/include" -query-api-mapping=tex2DLayered | FileCheck %s -check-prefix=tex2DLayered +// tex2DLayered: CUDA API: +// tex2DLayered-NEXT: tex2DLayered(tex /*cudaTextureObject_t*/, x /*float*/, y /*float*/, +// tex2DLayered-NEXT: layer /*int*/); +// tex2DLayered-NEXT: tex2DLayered(ptr /*T **/, tex /*cudaTextureObject_t*/, x /*float*/, +// tex2DLayered-NEXT: y /*float*/, layer /*int*/); +// tex2DLayered-NEXT: tex2DLayered(tex /*cudaTextureObject_t*/, x /*float*/, y /*float*/, +// tex2DLayered-NEXT: layer /*int*/, is_resident /*bool **/); +// tex2DLayered-NEXT: tex2DLayered(ptr /*T **/, tex /*cudaTextureObject_t*/, x /*float*/, +// tex2DLayered-NEXT: y /*float*/, layer /*int*/, is_resident /*bool **/); +// tex2DLayered-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// tex2DLayered-NEXT: sycl::ext::oneapi::experimental::sample_image_array(tex, sycl::float2(x, y), layer); +// tex2DLayered-NEXT: *ptr = sycl::ext::oneapi::experimental::sample_image_array(tex, sycl::float2(x, y), layer); +// tex2DLayered-NEXT: sycl::ext::oneapi::experimental::sample_image_array(tex, sycl::float2(x, y), is_resident); +// tex2DLayered-NEXT: *ptr = sycl::ext::oneapi::experimental::sample_image_array(tex, sycl::float2(x, y), is_resident); diff --git a/clang/test/dpct/query_api_mapping/Runtime/test-after101.cu b/clang/test/dpct/query_api_mapping/Runtime/test-after101.cu new file mode 100644 index 000000000000..b17c032d776d --- /dev/null +++ b/clang/test/dpct/query_api_mapping/Runtime/test-after101.cu @@ -0,0 +1,8 @@ +// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0 +// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0 + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cudaStreamBeginCapture | FileCheck %s -check-prefix=CUDASTREAMBEGINCAPTURE +// CUDASTREAMBEGINCAPTURE: CUDA API: +// CUDASTREAMBEGINCAPTURE-NEXT: cudaStreamBeginCapture(s /*cudaStream_t*/, sc /*cudaStreamCaptureMode*/); +// CUDASTREAMBEGINCAPTURE-NEXT: Is migrated to (with the option --use-experimental-features=graph): +// CUDASTREAMBEGINCAPTURE-NEXT: dpct::experimental::begin_recording(s); diff --git a/clang/test/dpct/query_api_mapping/Runtime/test-after102.cu b/clang/test/dpct/query_api_mapping/Runtime/test-after102.cu new file mode 100644 index 000000000000..47d42a31564e --- /dev/null +++ b/clang/test/dpct/query_api_mapping/Runtime/test-after102.cu @@ -0,0 +1,10 @@ +// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1 +// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1 + +// RUN: dpct --cuda-include-path="%cuda-path/include" -query-api-mapping=cudaGraphExecUpdate | FileCheck %s -check-prefix=cudaGraphExecUpdate +// cudaGraphExecUpdate: CUDA API: +// cudaGraphExecUpdate-NEXT: cudaGraphExecUpdate(graph_exec /*cudaGraphExec_t*/, graph /*cudaGraph_t*/, +// cudaGraphExecUpdate-NEXT: node /*cudaGraphNode_t **/, +// cudaGraphExecUpdate-NEXT: result /*cudaGraphExecUpdateResult **/); +// cudaGraphExecUpdate-NEXT: Is migrated to (with the option --use-experimental-features=graph): +// cudaGraphExecUpdate-NEXT: graph_exec->update(*graph);