Mma operation is recommanded for the calculation of D = A * B + C in fp16. However, many of the tutorial about how to use mma is just use shared memory to do everything. I wonder if there are some tutorials that use register to store C & D through the whole process?

I’m not going to use the word “coalesce” here, although I know what you mean, I think. You are interested in arguments being packed in a sensible fashion into registers. I would suggest (maybe others will disagree with me) that these questions cannot be fully answered or are not best answered by s…

The cuda programming guide only talks about how to fill in the PTX blanks by registers. e.g: .reg .f16x2 %Ra<4>, %Rb<2>, %Rc<2>, %Rd<2>; mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 {%Rd0, %Rd1}, {%Ra0, %Ra1, %Ra2, %Ra3}, {%Rb0, %Rb1}, {%Rc0, %Rc1}; However, as a register has 32 bits, I…

Should I use __half2 type to save the values in the register? e.g global f() { __half2 Rc[32]; __half2 Rd[32]; … }

[image] Error or incomprehension, MMa ptx mixed precision Bfloat16 rtx3080 CUDA Programming and Performance As an example demonstrating a “full” M=16, N =8, K=8 matrix multiply, let’s perform the following op: D = A*B + C we will set C to all zero. We will choose these …

Will compiler automatically coalesce the register in order to use less register? For example, in your code, you define “half a[4]”. Will it use 2 registers or 4 registers in the end?

Do I need to coalesce the registers by myself?

[image] half-0: For example, I allocate 32 registers per thread for Rc & Rd, they ought to save 64 fp16 values. m16n8k16 has C and D of matrix size m x n = 16 * 8 = 128 values each, not 64 values each. That is 2 32-bit registers with 2 16-bit values each over 32 threads (or lanes) per warp =…

Another question：If the calculation is C = A * B + C, can I just write code in this way? float const *C = reinterpret_cast<float const *>(&c); float *D = reinterpret_cast<float *>(&c);

I don’t think that will work, did you try it? Types matter, of course, and this includes when passing data into PTX registers. But you’ve not provided a complete code to demonstrate the change, so I can’t be certain of what you are suggesting. When PTX wants a u32 register, a floating point quant…

like this __global__ void mma_fp16_acc_fp32(float *out) { float c[4] = {0., 0., 0., 0.}; float d[4] = {0., 0., 0., 0.}; half a[4] = {1., 1., 1., 1.}; half b[2] = {1., 1.}; // the above would set our input matrices to all 1 // now lets modify some values if (threadIdx.x%4…

Use register for mma calculation results store

Accelerated Computing CUDA CUDA Programming and Performance

Robert_Crovella October 31, 2024, 3:32pm 4

Topic		Replies	Views
Why compiler don't use registers to store my data? CUDA Programming and Performance	42	746	December 7, 2024
Error or incomprehension, MMa ptx mixed precision Bfloat16 rtx3080 CUDA Programming and Performance	19	2764	April 12, 2021
Complete minimal ptx example for: mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 CUDA Programming and Performance	3	1250	October 25, 2024
Register in kernel CUDA Programming and Performance	3	107	November 17, 2024
About compute accuracy CUDA Programming and Performance	22	365	February 10, 2025
.reg f16x2 %Rb<1> in ISA example: mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 CUDA Programming and Performance	1	105	October 25, 2024
Padding of mma operation CUDA Programming and Performance	20	502	December 19, 2024
Wrong answer with mma.sync.aligned.m8n8k4 CUDA Programming and Performance cuda , kernel	7	1610	April 3, 2023
# of registers in different for different datatypes CUDA Programming and Performance	3	644	January 21, 2020
Problem with the instruction "mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32" CUDA Programming and Performance	2	2032	June 18, 2021

Use register for mma calculation results store

Related topics