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

Use syclcompat dim3 #97

Merged
merged 1 commit into from
Jul 15, 2024
Merged

Use syclcompat dim3 #97

Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
210 changes: 101 additions & 109 deletions include/cutlass/gpu_generics.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -294,128 +294,120 @@ CUTLASS_DEVICE T hfma2(const T a, const T b, const T c) {

namespace cutlass {

// Stream
using cudaStream_t = void *;
// Stream
using cudaStream_t = void *;

// dim3
struct dim3 {
uint x, y, z;
using dim3 = syclcompat::dim3;

dim3() = default;
// Atomic

dim3(uint x, uint y, uint z) : x(x), y(y), z(z) {}
};


// Atomic

CUTLASS_DEVICE int atomicAdd(int *address, int val) {
CUTLASS_DEVICE int atomicAdd(int *address, int val) {
#if defined(__SYCL_DEVICE_ONLY__)
return syclcompat::atomic_fetch_add(address, val);
return syclcompat::atomic_fetch_add(address, val);
#endif
return 0;
}
return 0;
}

CUTLASS_DEVICE int atomicCAS(int *address, int compare, int val) {
CUTLASS_DEVICE int atomicCAS(int *address, int compare, int val) {
#if defined(__SYCL_DEVICE_ONLY__)
syclcompat::atomic_compare_exchange_strong(address, compare, val);
syclcompat::atomic_compare_exchange_strong(address, compare, val);
#endif
return 0;
}

// Error
using cudaError_t = unsigned int;
constexpr cudaError_t cudaSuccess = 0;
constexpr cudaError_t cudaErrorUnknown = 100;

CUTLASS_HOST_DEVICE
const char *cudaGetErrorString(cudaError_t error) {
return "";
}

CUTLASS_HOST_DEVICE
void cuGetErrorString(cudaError_t error, const char **) {
}

CUTLASS_HOST
cudaError_t cudaGetLastError() {
return cudaSuccess;
}

CUTLASS_HOST_DEVICE
cudaError_t cudaGetDevice(int *device) {
return cudaSuccess;
}

// Mem copy
enum cudaMemcpyKind {
cudaMemcpyHostToHost = 0,
cudaMemcpyHostToDevice = 1,
cudaMemcpyDeviceToHost = 2,
cudaMemcpyDeviceToDevice = 3
};

CUTLASS_HOST_DEVICE
cudaError_t cudaMemsetAsync(void *devPtr, unsigned int value, size_t count, cudaStream_t stream = nullptr) {
syclcompat::fill_async(devPtr, value, count);
return cudaSuccess;
}

using CUresult = unsigned int;
using CUdeviceptr = unsigned int*;
constexpr CUresult CUDA_SUCCESS = 0;

CUTLASS_HOST_DEVICE
CUresult cuMemsetD32Async(CUdeviceptr devPtr, uint32_t value, size_t count, cudaStream_t stream = nullptr) {
void *ptr = reinterpret_cast<void *>(devPtr);
syclcompat::fill_async(ptr, value, count);
return cudaSuccess;
}

CUTLASS_HOST_DEVICE
CUresult cuMemsetD16Async(CUdeviceptr devPtr, uint16_t value, size_t count, cudaStream_t stream = nullptr) {
void *ptr = reinterpret_cast<void *>(devPtr);
syclcompat::fill_async(ptr, value, count);
return cudaSuccess;
}

CUTLASS_HOST_DEVICE
CUresult cuMemsetD8Async(CUdeviceptr devPtr, uint8_t value, size_t count, cudaStream_t stream = nullptr) {
void *ptr = reinterpret_cast<void *>(devPtr);
syclcompat::fill_async(ptr, value, count);
return cudaSuccess;
}

// FuncAttribute
using cudaFuncAttribute = unsigned int;
constexpr cudaFuncAttribute cudaFuncAttributeMaxDynamicSharedMemorySize = 0;

CUTLASS_HOST
cudaError_t cudaFuncSetAttribute(const void *func, cudaFuncAttribute attr, int value) {
return cudaSuccess;
}

using cudaDeviceAttr = unsigned int;
constexpr cudaDeviceAttr cudaDevAttrMultiProcessorCount = 0;

CUTLASS_HOST_DEVICE
cudaError_t cudaDeviceGetAttribute(int *value, cudaDeviceAttr attr, int device) {
return cudaSuccess;
}

constexpr unsigned int cudaOccupancyDisableCachingOverride = 0;

CUTLASS_HOST
cudaError_t cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
int *numBlocks, const void *func, int blockSize, size_t dynamicSMemSize, unsigned int flags) {
return cudaSuccess;
}
return 0;
}

// Error
using cudaError_t = unsigned int;
constexpr cudaError_t cudaSuccess = 0;
constexpr cudaError_t cudaErrorUnknown = 100;

CUTLASS_HOST_DEVICE
const char *cudaGetErrorString(cudaError_t error) {
return "";
}

CUTLASS_HOST_DEVICE
void cuGetErrorString(cudaError_t error, const char **) {
}

CUTLASS_HOST
cudaError_t cudaGetLastError() {
return cudaSuccess;
}

CUTLASS_HOST_DEVICE
cudaError_t cudaGetDevice(int *device) {
return cudaSuccess;
}

// Mem copy
enum cudaMemcpyKind {
cudaMemcpyHostToHost = 0,
cudaMemcpyHostToDevice = 1,
cudaMemcpyDeviceToHost = 2,
cudaMemcpyDeviceToDevice = 3
};

CUTLASS_HOST_DEVICE
cudaError_t cudaMemsetAsync(void *devPtr, unsigned int value, size_t count, cudaStream_t stream = nullptr) {
syclcompat::fill_async(devPtr, value, count);
return cudaSuccess;
}

using CUresult = unsigned int;
using CUdeviceptr = unsigned int*;
constexpr CUresult CUDA_SUCCESS = 0;

CUTLASS_HOST_DEVICE
CUresult cuMemsetD32Async(CUdeviceptr devPtr, uint32_t value, size_t count, cudaStream_t stream = nullptr) {
void *ptr = reinterpret_cast<void *>(devPtr);
syclcompat::fill_async(ptr, value, count);
return cudaSuccess;
}

CUTLASS_HOST_DEVICE
CUresult cuMemsetD16Async(CUdeviceptr devPtr, uint16_t value, size_t count, cudaStream_t stream = nullptr) {
void *ptr = reinterpret_cast<void *>(devPtr);
syclcompat::fill_async(ptr, value, count);
return cudaSuccess;
}

CUTLASS_HOST_DEVICE
CUresult cuMemsetD8Async(CUdeviceptr devPtr, uint8_t value, size_t count, cudaStream_t stream = nullptr) {
void *ptr = reinterpret_cast<void *>(devPtr);
syclcompat::fill_async(ptr, value, count);
return cudaSuccess;
}

// FuncAttribute
using cudaFuncAttribute = unsigned int;
constexpr cudaFuncAttribute cudaFuncAttributeMaxDynamicSharedMemorySize = 0;

CUTLASS_HOST
cudaError_t cudaFuncSetAttribute(const void *func, cudaFuncAttribute attr, int value) {
return cudaSuccess;
}

using cudaDeviceAttr = unsigned int;
constexpr cudaDeviceAttr cudaDevAttrMultiProcessorCount = 0;

CUTLASS_HOST_DEVICE
cudaError_t cudaDeviceGetAttribute(int *value, cudaDeviceAttr attr, int device) {
return cudaSuccess;
}

constexpr unsigned int cudaOccupancyDisableCachingOverride = 0;

CUTLASS_HOST
cudaError_t cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
int *numBlocks, const void *func, int blockSize, size_t dynamicSMemSize, unsigned int flags) {
return cudaSuccess;
}

} // cutlass namespace

// Expose dim3 in the cute namespace
namespace cute {
using dim3 = cutlass::dim3;
using dim3 = syclcompat::dim3;
}
#endif

Expand Down