Opcode 30、57、58 for DS

[Yzgll]

Teacher Sun, hello, I was recently researching the emulator's local memory, and I wanted to load some global memory data into local memory for use, so I added three new instructions. DS_WRITE_B8: DS[A] = D0[7:0]; byte write, DS_READ_I8: R = signext(DS[A][7:0]}; signed byte read and DS_READ_U8: R = {24 'h0,DS[A][7:0]}; unsigned byte read. However, after this I got the wrong result, which was found to be due to the inconsistency between the data loaded into local memory and the data in global memory. Therefore, I suspect that the three instruction functions added are not correct, but I have checked for several times but could not find the error, so I would like to ask you to help me check where the error is. Looking forward to your reply.

func (u *ALUImpl) runDSWRITEB8(state InstEmuState) {
	inst := state.Inst()
	sp := state.Scratchpad()
	layout := sp.AsDS()
	lds := u.LDS()

	i := uint(0)
	for i = 0; i < 64; i++ {
		if !laneMasked(layout.EXEC, i) {
			continue
		}

		addr := layout.ADDR[i] + inst.Offset0
		data0offset := uint(8 + 64*4)
		byteToWrite := sp[data0offset+i*16]
		lds[addr] = byteToWrite
	}
}

func (u *ALUImpl) runDSREADI8(state InstEmuState) {
	inst := state.Inst()
	sp := state.Scratchpad()
	layout := sp.AsDS()
	lds := u.LDS()

	i := uint(0)
	for i = 0; i < 64; i++ {
		if !laneMasked(layout.EXEC, i) {
			continue
		}

		addr0 := layout.ADDR[i] + inst.Offset0
		bytedata := lds[addr0]
		signedData := int32(int8(bytedata)&0xFF)
		dstOffset := uint(8 + 64*4 + 256*4*2)
		sp[dstOffset] = byte(signedData)
		sp[dstOffset+1] = byte(signedData >> 8)
		sp[dstOffset+2] = byte(signedData >> 16)
		sp[dstOffset+3] = byte(signedData >> 24)

	}
}

func (u *ALUImpl) runDSREADU8(state InstEmuState) {
	inst := state.Inst()
	sp := state.Scratchpad()
	layout := sp.AsDS()
	lds := u.LDS()
	i := uint(0)
	for i = 0; i < 64; i++ {
		if !laneMasked(layout.EXEC, i) {
			continue
		}
		addr0 := layout.ADDR[i] + inst.Offset0
		byteData := lds[addr0]
		usignedData := uint32(uint8(byteData)0xFF)
		dstOffset := uint(8 + 64*4 + 256*4*2)
		sp[dstOffset] = byte(usignedData)
		sp[dstOffset+1] = byte(usignedData >> 8)
		sp[dstOffset+2] = byte(usignedData >> 16)
		sp[dstOffset+3] = byte(usignedData >> 24)

	}
}

[syifan]

I do not understand how you added new instructions. How can you make the AMD compilers generate the instructions? Can I see the assembly?

[Yzgll]

I just added case 30, case 57, and case 58 to aluds.go, and implemented the corresponding functions runDSWRITEB8(state), runDSREADI8(state), and runDSREADU8(state) below to simulate the instruction. And when you mentioned "assembly", were you referring to kernels.hsaco?

[syifan]

Or what is your OpenCL program? How do you know the added instructions are used in the hsaco file?

[Yzgll]

Here is the kernel of my OpenCL program, with comments added where I made modifications. After modifying the kernel and before adding the instructions, the program panicked and could not run, displaying the message: "Opcode 30, 57 for DS format is not implemented". After I added the instructions, the program was able to run successfully, but the computation results were incorrect. Therefore, it seems to me that these instructions are used in the hsaco file, and the computation errors are due to mistakes in the code I wrote.

__kernel void Encrypt(__global uchar* input, 
                      __global uint* expanded_key, 
                      __global uchar* s) {
  uchar state[16];
  int tid = get_global_id(0);
  
  //load s_box to local memory
  __local uchar local_s[256];
  for(int i = 0;i < 256;i ++){
    local_s[i]=s[i];
  }

  for (int i = 0; i < 16; i++) {
    state[i] = input[tid * 16 + i];
  }

  AddRoundKey(state, expanded_key, 0);
  for (int i = 1; i < 10; i++) {
    SubBytes(state, local_s);//use local_s not s
    ShiftRows(state);
    MixColumns(state);
    AddRoundKey(state, expanded_key, i * 4);
  }
  SubBytes(state, local_s);
  ShiftRows(state);
  AddRoundKey(state, expanded_key, 10 * 4);

  for (int i = 0; i < 16; i++) {
    input[tid * 16 + i] = state[i];
  }
}

[Yzgll]

Or what is your OpenCL program? How do you know the added instructions are used in the hsaco file?

You are right. I set the input as fixed and unchangeable. No matter how I modify the three added instruction functions, the resulting output is still incorrect and does not change with the modifications to the instruction functions. I even set the added instruction functions as empty functions and still got the same error result. It seems that the instruction functions I added only allow the program to run but do not implement the corresponding functionality. Why is this happening?

[syifan]

Do not use the stub function. Use the logWholeWF instead.

Test files are for tests, and they are not part of the MGPUSim simulations. If you want to run tests, you can run go test in the folder that contains the test.

[Yzgll]

As you said, I uncommented and only used logWholeWF, but still nothing on the browser side. I just followed the steps in readme to rename the debug file and start the http server. The second problem is that go test can not run the test file, showing something undefined, why?

[syifan]

Can you try removing the two if statements and just keep the h.logWholeWf?

[Yzgll]

Can I interpret it as a direct call to h.logWholeWf(wf)?

[syifan]

Yes.