444 commits

This branch

This branch

All branches

Author	SHA1	Message	Date
Andrzej Janik	9a3ac6e36a	Update docs/src/supported-sw.md ... Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>	2025-08-26 18:17:02 -07:00
Andrzej Janik	58762d62ba	Update docs/src/usage.md ... Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>	2025-08-26 18:16:53 -07:00
Andrzej Janik	a45993ce0b	Improve wording add tracing section	2025-08-26 23:10:08 +00:00
Andrzej Janik	6c25e32412	Fix mdbook path again	2025-08-14 18:04:32 -07:00
Andrzej Janik	cdc05b631c	Fix mdbook path	2025-08-14 17:55:31 -07:00
Andrzej Janik	f82c327c52	Update docs	2025-08-14 17:49:03 -07:00
Andrzej Janik	5668a81111	Update mdbook path	2025-08-14 21:30:55 +00:00
Andrzej Janik	8f76c3ab57	change mdbook paths	2025-08-14 19:59:45 +00:00
Andrzej Janik	c761ea92b6	Try use both dimensions for icons	2025-08-14 19:49:17 +00:00
Andrzej Janik	88958e604b	Update main page layout	2025-08-14 18:50:12 +00:00
Andrzej Janik	748b2e5e28	Try center	2025-08-14 18:18:39 +00:00
Andrzej Janik	ad105f6971	Start working on docs	2025-08-14 18:15:55 +00:00
Violet	a420601128	Add test for unrecognized statement error with vector braces (#472) ... The old code using `take_till_inclusive` assumed that a right brace would be the end of a block and therefore never part of a statement. However, some PTX statements can include vector operands. This meant that any unrecognized statement with a vector operand would backtrace and eventually produce an unhelpful context error rather than an `UnrecognizedStatement` error. This pull request also adds a mechanism for testing parser errors.	2025-08-13 17:23:51 -07:00
Joëlle van Essen	fe7a18f912	zoc (ZLUDA offline compiler) (#344)	2025-08-13 15:27:02 -07:00
Andrzej Janik	28eca3d75a	Implement kernel cache (#465)	2025-08-11 09:20:40 -07:00
Violet	d2f92e4267	More tracing for custom parsers (#471)	2025-08-08 12:19:27 -07:00
Violet	ef98c1e0ba	Add tracing to custom parsers (#469)	2025-08-07 11:51:56 -07:00
Violet	94eec34bdb	More descriptive syntax errors (#466) ... * More descriptive syntax errors * cargo fmt * Remove brackets	2025-08-05 17:23:55 -07:00
Andrzej Janik	a1b7600718	Fix version in nightly trigger (#464)	2025-08-04 13:09:37 -07:00
Violet	dd05752fc4	Implement shf instruction (#463) ... * Implement `shf` instruction * Tests for `shf` * cargo fmt	2025-08-04 10:15:46 -07:00
Andrzej Janik	ce9c0aac23	In nightly runs, make sure cargo is in $PATH (#462)	2025-08-01 09:36:11 -07:00
Andrzej Janik	c00496b92a	Install curl correctly (#461)	2025-07-31 19:56:37 -07:00
Andrzej Janik	7fc6caffb6	Update nightly runner's ROCm version to 6.3.4 (#460)	2025-07-31 19:46:49 -07:00
Andrzej Janik	52d3ea624c	Fix nightly run (#459)	2025-07-31 19:20:39 -07:00
Andrzej Janik	704a94e6f4	Fix nightly tests trigger (#458)	2025-07-31 18:31:54 -07:00
Andrzej Janik	cd7e2f8e36	Force loading ZLUDA through LD_PRELOAD (#447) ... Certain applications (pytorch) decide that it's a great idea to distribute whole CUDA driver and link to it with DT_RPATH. This igores LD_LIBRARY_PATH. This code defeats that evil mechanism through any means necessary	2025-07-31 18:00:13 -07:00
Violet	96ae27e9e1	Implement cublas functions needed for llm.c (#457)	2025-07-31 11:08:53 -07:00
Violet	99c36092be	Use FromCuda in zluda_blas (#455)	2025-07-31 09:52:10 -07:00
Andrzej Janik	49aabffdcc	Rename zluda_dump to zluda_trace (#456) ... * Rename zluda_dump to zluda_trace * Minor naming fixes	2025-07-31 08:07:03 -07:00
Violet	4d163a4d9b	Implement cuModuleGetGlobal_v2 (#454)	2025-07-30 16:34:21 -07:00
Violet	66db19a061	Move FromCuda and ZludaObject into a common crate (#452) ... * Refactor FromCuda error type to be generic * Create zluda_common crate * Move FromCuda trait into zluda_common * Write some doc comments * Fix typo * Edit comment * Fix formatting	2025-07-30 15:53:22 -07:00
Violet	b8bcbec295	Always use Unix line endings (#453)	2025-07-30 15:09:47 -07:00
Violet	21ef5f60a3	Check Rust formatting on pull requests (#451) ... * Check Rust formatting on pull requests This should help us maintain consistent style, without having unrelated style changes in pull requests from running `rustfmt`. * cargo fmt non-generated files * Ignore generated files	2025-07-30 14:55:09 -07:00
Violet	98b601d15a	Use normalize_fn for performance libraries (#449) ... The goal here is to make the performance library implementations work more like zluda.	2025-07-30 14:02:01 -07:00
Violet	c07d7678cd	Format files (#450)	2025-07-30 10:30:49 -07:00
Violet	481c3550fa	Convert CUDA performance lib statuses to Rust result types (#444) ... These changes replicate how the main library is handled. cuDNN still needs to have zluda_bindgen run and zluda_dump_common updated	2025-07-29 14:28:14 -07:00
Violet	303e4c2fb2	Update rocm_setup_build.sh (#446)	2025-07-29 11:05:16 -07:00
Violet	4ffa669cce	Fix Windows linkage (#445)	2025-07-29 10:15:32 -07:00
Violet	d81404eb70	Add support for rocblas to zluda_bindgen (#440) ... One step of several for adding cublas support	2025-07-28 15:07:22 -07:00
Andrzej Janik	8dbc7208de	Try to make ZLUDA more robust on Windows (#442) ... On my machine ZLUDA seems to segfault when initializing LLVM's C++ statics in Blender. Blender ships with C++ runtime. It seems that compiling C++ runtime statically fixes the issue. Might be actually unrelated. Additionally, dtor crate on Windows seem to use a slightly dodgy method, so replace it with something more straightforward	2025-07-28 13:20:04 -07:00
Violet	f192dd317a	Use implicit FromCuda for library::get_module (#439)	2025-07-28 06:42:14 -07:00
Violet	8c23ef1ded	Rename cuda_base cuda_macros (#435)	2025-07-25 11:09:50 -07:00
Violet	ba38da0bbc	Silence unused variable warnings (#434) ... I'd left these in originally because I'd assumed that these functions would need full implementations soon, but they're really annoying. I've fixed all the other compiler warnings as well.	2025-07-25 10:28:33 -07:00
Andrzej Janik	c1dda55235	Add nightly tests (#433)	2025-07-24 16:14:06 -07:00
Andrzej Janik	e8e20294a6	Set newly created context as current (#431)	2025-07-24 11:05:04 -07:00
Andrzej Janik	5deada8426	Add cuCtxCreate_v2 and cuCtxDestroy_v2 (#430)	2025-07-23 17:33:59 -07:00
Violet	2b90fdb56c	Add support for cp.async (#427) ... Adds support for * `cp.async` * `cp.async.commit_group` * `cp.async.wait_group` * `cp.async.wait_all` Asynchronous copy operations are only supported by AMD Instinct GPUs, so for now we lower them as synchronous copy operations. Because of this, `cp.async.commit_group`, `cp.async.wait_group`, and `cp.async.wait_all` are no-op.	2025-07-23 16:25:49 -07:00
Andrzej Janik	3746079b1a	Assorted instruction fixes (#423) ... This fixes transcendentals and some other buggy instructions exposed by `ptx_tests` (abs, neg). Add (slow - hardware limitation) tanh. Only two remaining incorrect instructions are div and sqrt with non-default rounding, but this commit is already bloated enough	2025-07-23 15:50:35 -07:00
Violet	119b635b9d	Emit correct alignment for loads and stores (#429)	2025-07-23 14:55:52 -07:00
Violet	a86ba3d642	Remove Type::Pointer (#428)	2025-07-23 11:22:17 -07:00