Skip to content

Updating SYCL For CUDA examples to CUDA 12 and latest intel llvm tip #27

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 7 commits into from
Apr 19, 2023
Merged

Updating SYCL For CUDA examples to CUDA 12 and latest intel llvm tip #27

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
245ccd0
Ignoring VIM temporary files
Ruyk Feb 9, 2023
0c886cb
Vector Addition example update
Ruyk Feb 9, 2023
283e69d
Removing unnecessary CUDA Driver types
Ruyk Feb 13, 2023
0d4facb
Using moder queue construction
Ruyk Feb 13, 2023
6cfee58
Format files
Ruyk Feb 13, 2023
ce4abc2
Explicitly setting CUDA context on host task
Ruyk Apr 19, 2023
a06022b
Addressing feedback from Gordon
Ruyk Apr 19, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 4 additions & 1 deletion .gitignore
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,4 +34,7 @@
# Temporaries
*~
*#
*/build
*/build
]
# vim
*.swp
38 changes: 12 additions & 26 deletions examples/sgemm_interop/sycl_sgemm.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,8 +25,7 @@
#include <iostream>
#include <vector>

#include <CL/sycl.hpp>
#include <CL/sycl/backend/cuda.hpp>
#include <sycl/sycl.hpp>

#include <cublas_v2.h>
#include <cuda.h>
Expand All @@ -47,25 +46,6 @@ void inline checkCudaErrorMsg(cudaError status, const char *msg) {
}
}

void inline checkCudaErrorMsg(CUresult status, const char *msg) {
if (status != CUDA_SUCCESS) {
std::cout << "ERROR CUDA: " << msg << " - " << status << std::endl;
exit(EXIT_FAILURE);
}
}

class CUDASelector : public sycl::device_selector {
public:
int operator()(const sycl::device &device) const override {
if(device.get_platform().get_backend() == sycl::backend::ext_oneapi_cuda){
std::cout << " CUDA device found " << std::endl;
return 1;
} else{
return -1;
}
}
};

int main() {
using namespace sycl;

Expand All @@ -88,7 +68,9 @@ int main() {
// B is a matrix fill with 1
std::fill(std::begin(h_B), std::end(h_B), 1.0f);

sycl::queue q{CUDASelector()};
sycl::queue q{[](auto &d) {
return (d.get_platform().get_backend() == sycl::backend::ext_oneapi_cuda);
}};

cublasHandle_t handle;
CHECK_ERROR(cublasCreate(&handle));
Expand All @@ -104,12 +86,16 @@ int main() {
auto d_C = b_C.get_access<sycl::access::mode::write>(h);

h.host_task([=](sycl::interop_handle ih) {
cuCtxSetCurrent(ih.get_native_context<backend::ext_oneapi_cuda>());
// Set the correct cuda context & stream
cuCtxSetCurrent(ih.get_native_context<backend::ext_oneapi_cuda>());
auto cuStream = ih.get_native_queue<backend::ext_oneapi_cuda>();
cublasSetStream(handle, cuStream);
auto cuA = reinterpret_cast<float *>(ih.get_native_mem<backend::ext_oneapi_cuda>(d_A));
auto cuB = reinterpret_cast<float *>(ih.get_native_mem<backend::ext_oneapi_cuda>(d_B));
auto cuC = reinterpret_cast<float *>(ih.get_native_mem<backend::ext_oneapi_cuda>(d_C));
auto cuA = reinterpret_cast<float *>(
ih.get_native_mem<backend::ext_oneapi_cuda>(d_A));
auto cuB = reinterpret_cast<float *>(
ih.get_native_mem<backend::ext_oneapi_cuda>(d_B));
auto cuC = reinterpret_cast<float *>(
ih.get_native_mem<backend::ext_oneapi_cuda>(d_C));

CHECK_ERROR(cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, WIDTH, HEIGHT,
WIDTH, &ALPHA, cuA, WIDTH, cuB, WIDTH, &BETA,
Expand Down
60 changes: 20 additions & 40 deletions examples/sgemm_interop/sycl_sgemm_usm.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,8 +25,7 @@
#include <iostream>
#include <vector>

#include <CL/sycl.hpp>
#include <CL/sycl/backend/cuda.hpp>
#include <sycl/sycl.hpp>

#include <cublas_v2.h>
#include <cuda.h>
Expand All @@ -47,25 +46,6 @@ void inline checkCudaErrorMsg(cudaError status, const char *msg) {
}
}

void inline checkCudaErrorMsg(CUresult status, const char *msg) {
if (status != CUDA_SUCCESS) {
std::cout << "ERROR CUDA: " << msg << " - " << status << std::endl;
exit(EXIT_FAILURE);
}
}

class CUDASelector : public sycl::device_selector {
public:
int operator()(const sycl::device &device) const override {
if(device.get_platform().get_backend() == sycl::backend::ext_oneapi_cuda){
std::cout << " CUDA device found " << std::endl;
return 1;
} else{
return -1;
}
}
};

int main() {
using namespace sycl;

Expand All @@ -88,12 +68,14 @@ int main() {
// B is a matrix fill with 1
std::fill(std::begin(h_B), std::end(h_B), 1.0f);

sycl::queue q{CUDASelector()};
sycl::queue q{[](auto &d) {
return (d.get_platform().get_backend() == sycl::backend::ext_oneapi_cuda);
}};

// Allocate memory on the device
float* d_A = sycl::malloc_device<float>(WIDTH*HEIGHT,q);
float* d_B = sycl::malloc_device<float>(WIDTH*HEIGHT,q);
float* d_C = sycl::malloc_device<float>(WIDTH*HEIGHT,q);
float *d_A = sycl::malloc_device<float>(WIDTH * HEIGHT, q);
float *d_B = sycl::malloc_device<float>(WIDTH * HEIGHT, q);
float *d_C = sycl::malloc_device<float>(WIDTH * HEIGHT, q);

// Copy matrices A & B to device from host vectors
const size_t numBytes = WIDTH * HEIGHT * sizeof(float);
Expand All @@ -105,21 +87,19 @@ int main() {
CHECK_ERROR(cublasCreate(&handle));

q.submit([&](handler &h) {

h.host_task([=](sycl::interop_handle ih) {

// Set the correct cuda context & stream
cuCtxSetCurrent(ih.get_native_context<backend::ext_oneapi_cuda>());
auto cuStream = ih.get_native_queue<backend::ext_oneapi_cuda>();
cublasSetStream(handle, cuStream);

// Call generalised matrix-matrix multiply
CHECK_ERROR(cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, WIDTH, HEIGHT,
WIDTH, &ALPHA, d_A, WIDTH, d_B, WIDTH, &BETA,
d_C, WIDTH));
cuStreamSynchronize(cuStream);
});
}).wait();
h.host_task([=](sycl::interop_handle ih) {
// Set the correct cuda context & stream
cuCtxSetCurrent(ih.get_native_context<backend::ext_oneapi_cuda>());
auto cuStream = ih.get_native_queue<backend::ext_oneapi_cuda>();
cublasSetStream(handle, cuStream);

// Call generalised matrix-matrix multiply
CHECK_ERROR(cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, WIDTH, HEIGHT,
WIDTH, &ALPHA, d_A, WIDTH, d_B, WIDTH, &BETA,
d_C, WIDTH));
cuStreamSynchronize(cuStream);
});
}).wait();

// Copy the result back to host
q.memcpy(h_C.data(), d_C, numBytes).wait();
Expand Down
3 changes: 0 additions & 3 deletions examples/vector_addition/README.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,9 +7,6 @@ to highlight how to build an application with SYCL for CUDA using DPC++ support,
for which an example CMakefile is provided. For detailed documentation on how to
migrate from CUDA to SYCL, see [SYCL For CUDA Developers](https://developer.codeplay.com/products/computecpp/ce/guides/sycl-for-cuda-developers).

Note currently the CUDA backend does not support the [USM](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/USM/USM.adoc) extension, so we use
`sycl::buffer` and `sycl::accessors` instead.

Pre-requisites
---------------

Expand Down
12 changes: 5 additions & 7 deletions examples/vector_addition/vector_addition.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,10 +34,9 @@ int main(int argc, char *argv[]) {

// Initialize input data
{
const auto dwrite_t = sycl::access::mode::discard_write;
sycl::host_accessor h_a{bufA, sycl::write_only};
sycl::host_accessor h_b{bufB, sycl::write_only};

auto h_a = bufA.get_access<dwrite_t>();
auto h_b = bufB.get_access<dwrite_t>();
for (int i = 0; i < N; i++) {
h_a[i] = sin(i) * sin(i);
h_b[i] = cos(i) * cos(i);
Expand All @@ -63,15 +62,14 @@ int main(int argc, char *argv[]) {
auto b = bufB.get_access<read_t>(h);
auto c = bufC.get_access<write_t>(h);

h.parallel_for(VecSize,
[=](sycl::id<1> i) { c[i] = a[i] + b[i]; });
h.parallel_for(VecSize, [=](sycl::id<1> i) { c[i] = a[i] + b[i]; });
};

myQueue.submit(cg);

{
const auto read_t = sycl::access::mode::read;
auto h_c = bufC.get_access<read_t>();
sycl::host_accessor h_c{bufC, sycl::read_only};

double sum = 0.0f;
for (int i = 0; i < N; i++) {
sum += h_c[i];
Expand Down