Update the mask on the signal event ID

b-sumner · b-sumner · commit 009964308929 · 2025-01-22T12:43:52.000-08:00
diff --git a/amd/device-libs/ockl/src/hsaqs.cl b/amd/device-libs/ockl/src/hsaqs.cl
@@ -72,7 +72,11 @@ update_mbox(const __global amd_signal_t *sig)
     if (mb) {
         uint id = sig->event_id;
         AS(mb, id, memory_order_release, memory_scope_all_svm_devices);
-        __builtin_amdgcn_s_sendmsg(1 | (0 << 4), __builtin_amdgcn_readfirstlane(id) & 0xff);
+        uint mid = id &
+                    (__oclc_ISA_version < 9000 ? 0xff :
+                     (__oclc_ISA_version < 10000 ? 0xffffff :
+                      (__oclc_ISA_version < 11000 ? 0x7fffff : 0xffffff)));
+        __builtin_amdgcn_s_sendmsg(1 | (0 << 4), __builtin_amdgcn_readfirstlane(mid));
     }
 }
 

Original file line number	Diff line number	Diff line change
`@@ -72,7 +72,11 @@ update_mbox(const __global amd_signal_t *sig)`
`72`	`72`	`if (mb) {`
`73`	`73`	`uint id = sig->event_id;`
`74`	`74`	`AS(mb, id, memory_order_release, memory_scope_all_svm_devices);`
`75`		`- __builtin_amdgcn_s_sendmsg(1 \| (0 << 4), __builtin_amdgcn_readfirstlane(id) & 0xff);`
	`75`	`+ uint mid = id &`
	`76`	`+ (__oclc_ISA_version < 9000 ? 0xff :`
	`77`	`+ (__oclc_ISA_version < 10000 ? 0xffffff :`
	`78`	`+ (__oclc_ISA_version < 11000 ? 0x7fffff : 0xffffff)));`
	`79`	`+ __builtin_amdgcn_s_sendmsg(1 \| (0 << 4), __builtin_amdgcn_readfirstlane(mid));`
`76`	`80`	`}`
`77`	`81`	`}`
`78`	`82`