Earlier quoted context omitted.
Open AI™ (with a space)
I hope you're reading this Sam Altman: Make Open AI open . Or else you'll lose to the ecosystem.
Sam is busy with his new kiddo
DeepSeek open source DeepEP – library for MoE training and Inference
21–30 of 77 posts
Re: DeepSeek open source DeepEP – library for MoE training and Inference
#22> For extreme performance, we discover and use an out-of-doc PTX instruction: ld.global.nc.L1::no_allocate.L2::256B. This instruction will lead to an undefined behavior: accessing volatile GPU memory with non-coherent read-only PTX modifiers .nc. But the correctness is tested to be guaranteed with .L1::no_allocate on Hopper architectures, and performance will be much better.
Re: DeepSeek open source DeepEP – library for MoE training and Inference
#23The PTX instructions they talked about in the tech report should be pointing to the code here?
"For extreme performance, we discover and use a behavior-out-of-doc PTX instruction: ld.global.nc.L1::no_allocate.L2::256B. This instruction will lead to an undefined behavior: accessing volatile GPU memory with non-coherent read-only PTX modifiers .nc. But the correctness is tested to be guaranteed with .L1::no_allocate on Hopper architectures, and performance will be much better. If you find kernels not working on…
So non-coherent refers to bypassing cache coherency, ie don't care about what other units might have written to that address? And the L1/L2 modifiers are to avoid L1 thrashing, keeping the value in L2 only?
Or did I get that wrong?
Re: DeepSeek open source DeepEP – library for MoE training and Inference
#24Earlier quoted context omitted.
Open AI™ (with a space)
I hope you're reading this Sam Altman: Make Open AI open . Or else you'll lose to the ecosystem.
Too late, there is no more innovation from openai all the people that were the drivers left for Anthropic and the others. They had some of the biggest funding, had the advance... And yet they lost it.
Re: DeepSeek open source DeepEP – library for MoE training and Inference
#25Earlier quoted context omitted.
Open AI™ (with a space)
I hope you're reading this Sam Altman: Make Open AI open . Or else you'll lose to the ecosystem.
That’s an impossible ask. Sam is the pinnacle of capitalist ruling class, he’s a pure businessman. He has no interest in giving anything for free unless there’s a business plan. He doesn’t care about humanity. He’ll pretend to change the world and tell you that they’re inventing AGI, Q*, strawberry or whatever they’re branding it, but the reality is he knows it’s all over and unless there’s a major breakthrough this company will be in major financial trouble. Sorry for the rant but he doesn’t deserve much respect for turning all this science to grift. He’s actually the person the old openai board warned everyone about.
Re: DeepSeek open source DeepEP – library for MoE training and Inference
#26Is the PTX that everyone was looking forward to included this time?
> the PTX that everyone was looking forward to explanation for the rest of us why this is so important?
Parallel Thread Execution. Think of them as opcodes for the Nvidia GPUs.
They are a bit more complex that your traditional opcodes (the lowest level of abstraction accessible to users) in CPUs, as you can specify cache parameters, memory barriers etc.
There are documented combinations of parameters for those instructions but if you fuzz (search new combinations in a random or organized way because you hope some will work the way you want) you can find new ones with unexpected effects or with advantages (in various ways like not polluting caches, speed...)
Which is the case for example for ld.global.nc.L1::no_allocate.L2::256B that they use in deepseek that provides significant acceleration while beeing reliable (although not working on all architectures so they have ways to disable it)
Re: DeepSeek open source DeepEP – library for MoE training and Inference
#27Earlier quoted context omitted.
"For extreme performance, we discover and use a behavior-out-of-doc PTX instruction: ld.global.nc.L1::no_allocate.L2::256B. This instruction will lead to an undefined behavior: accessing volatile GPU memory with non-coherent read-only PTX modifiers .nc. But the correctness is tested to be guaranteed with .L1::no_allocate on Hopper architectures, and performance will be much better. If you find kernels not working on…
So non-coherent refers to bypassing cache coherency, ie don't care about what other units might have written to that address? And the L1/L2 modifiers are to avoid L1 thrashing, keeping the value in L2 only? Or did I get that wrong?
My understanding of the L2 part is that it asks for a 256b pre-fetch (only available on some platforms it seems) but they use vectors of 4 32bits signed ints max so not sure why only the 256 would work or if the fact that it did fetch the next 128 helps.
Re: DeepSeek open source DeepEP – library for MoE training and Inference
#28Earlier quoted context omitted.
Kind of ironic that DeepSeek is more Open than ChatGPT
Now. It’s amazing to me that everyone is like fuck OpenAI deepseek is the savior, when OpenAI’s papers and code jump started an AI revolution just a few years ago. Let’s wait the same number of years and see what deepseek does.
I thought the papers that jump started the revolution came from Google?
Re: DeepSeek open source DeepEP – library for MoE training and Inference
#29> For extreme performance, we discover and use an out-of-doc PTX instruction: ld.global.nc.L1::no_allocate.L2::256B. This instruction will lead to an undefined behavior: accessing volatile GPU memory with non-coherent read-only PTX modifiers .nc. But the correctness is tested to be guaranteed with .L1::no_allocate on Hopper architectures, and performance will be much better.
Practically speaking, is it possible for NVIDIA to "pull the rug" later, intentionally or otherwise, by subtly changing the behaviour of this out-of-doc instruction on new architectures?
Re: DeepSeek open source DeepEP – library for MoE training and Inference
#30Earlier quoted context omitted.
Now. It’s amazing to me that everyone is like fuck OpenAI deepseek is the savior, when OpenAI’s papers and code jump started an AI revolution just a few years ago. Let’s wait the same number of years and see what deepseek does.
I thought the papers that jump started the revolution came from Google?
Hinton. And if you'd ask himself probably Schmidthuber.