DeepSeek/DeepGEMM
1
0
Fork 0
mirror of https://github.com/deepseek-ai/DeepGEMM synced 2025-06-26 23:15:49 +00:00
Code Issues Actions Packages Projects Releases Wiki Activity

Merge pull request #80 from abcdabcd987/fix-link-error

Browse Source 
Fix linking error from ODR violation
This commit is contained in:
Chenggang Zhao 2025-04-07 09:31:58 +08:00 committed by GitHub
commit 3ea3cb203c
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
2 changed files with 4 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
deep_gemm/include/deep_gemm/mma_utils.cuh
View File

@ -888,15 +888,15 @@ struct SM90_U32x4_STSM_N {
} }
}; };
__device__ void warpgroup_arrive() { __forceinline__ __device__ void warpgroup_arrive() {
asm volatile("wgmma.fence.sync.aligned;\n" ::: "memory"); asm volatile("wgmma.fence.sync.aligned;\n" ::: "memory");
} }
__device__ void warpgroup_commit_batch() { __forceinline__ __device__ void warpgroup_commit_batch() {
asm volatile("wgmma.commit_group.sync.aligned;\n" ::: "memory"); asm volatile("wgmma.commit_group.sync.aligned;\n" ::: "memory");
} }
__device__ void warpgroup_fence_operand(float& reg) { __forceinline__ __device__ void warpgroup_fence_operand(float& reg) {
asm volatile("" : "+f"(reg) :: "memory"); asm volatile("" : "+f"(reg) :: "memory");
} }

2
deep_gemm/include/deep_gemm/tma_utils.cuh
View File

@ -40,7 +40,7 @@ constexpr CUtensorMapDataType get_CUtensorMapDataType() {
} }
} }
PFN_cuTensorMapEncodeTiled get_cuTensorMapEncodeTiled() { inline PFN_cuTensorMapEncodeTiled get_cuTensorMapEncodeTiled() {
// Get pointer to `cuTensorMapEncodeTiled` // Get pointer to `cuTensorMapEncodeTiled`
cudaDriverEntryPointQueryResult driver_status; cudaDriverEntryPointQueryResult driver_status;
void* cuTensorMapEncodeTiled_ptr = nullptr; void* cuTensorMapEncodeTiled_ptr = nullptr;