Skip to content
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

Fixed block dims and grid dims through syclcompat #50

Merged
merged 2 commits into from
Apr 29, 2024
Merged

Fixed block dims and grid dims through syclcompat #50

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
73c2db6
Fix block dims and grid dims
muhammad-tanvir-1211 Apr 26, 2024
0a14ac2
Fix call to thread0()
muhammad-tanvir-1211 Apr 29, 2024
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
14 changes: 8 additions & 6 deletions include/cutlass/cutlass.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -178,7 +178,7 @@ CUTLASS_HOST_DEVICE uint BlockDimX() {
#if defined(__CUDA_ARCH__)
return blockDim.x;
#elif defined(__SYCL_DEVICE_ONLY__)
return syclcompat::work_group_range::x();
return syclcompat::local_range::x();
#else
return 0;
#endif
Expand All @@ -188,7 +188,7 @@ CUTLASS_HOST_DEVICE uint BlockDimY() {
#if defined(__CUDA_ARCH__)
return blockDim.y;
#elif defined(__SYCL_DEVICE_ONLY__)
return syclcompat::work_group_range::y();
return syclcompat::local_range::y();
#else
return 0;
#endif
Expand All @@ -198,7 +198,7 @@ CUTLASS_HOST_DEVICE uint BlockDimZ() {
#if defined(__CUDA_ARCH__)
return blockDim.z;
#elif defined(__SYCL_DEVICE_ONLY__)
return syclcompat::work_group_range::z();
return syclcompat::local_range::z();
#else
return 0;
#endif
Expand All @@ -208,7 +208,7 @@ CUTLASS_HOST_DEVICE uint GridDimX() {
#if defined(__CUDA_ARCH__)
return gridDim.x;
#elif defined(__SYCL_DEVICE_ONLY__)
return syclcompat::global_range::x();
return syclcompat::work_group_range::x();
#else
return 0;
#endif
Expand All @@ -218,7 +218,7 @@ CUTLASS_HOST_DEVICE uint GridDimY() {
#if defined(__CUDA_ARCH__)
return gridDim.y;
#elif defined(__SYCL_DEVICE_ONLY__)
return syclcompat::global_range::y();
return syclcompat::work_group_range::y();
#else
return 0;
#endif
Expand All @@ -228,7 +228,7 @@ CUTLASS_HOST_DEVICE uint GridDimZ() {
#if defined(__CUDA_ARCH__)
return gridDim.z;
#elif defined(__SYCL_DEVICE_ONLY__)
return syclcompat::global_range::z();
return syclcompat::work_group_range::z();
#else
return 0;
#endif
Expand Down Expand Up @@ -372,6 +372,8 @@ CUTLASS_DEVICE int atomicCAS(int *address, int compare, int val) {
CUTLASS_HOST_DEVICE bool thread0() {
#if defined(__CUDA_ARCH__)
return (!threadIdx.x && !threadIdx.y && !threadIdx.z) && (!blockIdx.x && !blockIdx.y && !blockIdx.z);
#elif defined(CUTLASS_ENABLE_SYCL)
return (!syclcompat::global_id::x() && !syclcompat::global_id::y() && !syclcompat::global_id::z());
#else
return false;
#endif
Expand Down